From 4039c7d1af1e779d6a701cedec4025ce2b58ac87 Mon Sep 17 00:00:00 2001 From: eddierichter-amd Date: Thu, 25 Apr 2024 10:11:10 -0600 Subject: [PATCH] Cleaning up VCK5000 makefiles and removing test_vck5000.cpp and run_vck5000.lit for tests hat don't have the appropriate kernels (#1407) Co-authored-by: Joseph Melber --- .../basic/matrix_scalar_add/Makefile | 21 +-- .../basic/matrix_scalar_add/run_vck5000.lit | 7 +- .../basic/passthrough_dmas/Makefile | 9 +- .../basic/passthrough_dmas/run_vck5000.lit | 6 +- .../basic/vector_reduce_add/Makefile | 23 --- .../basic/vector_reduce_add/README.md | 17 -- .../basic/vector_reduce_add/run_vck5000.lit | 11 -- .../basic/vector_reduce_add/test_vck5000.cpp | 148 ------------------ .../basic/vector_reduce_max/Makefile | 21 --- .../basic/vector_reduce_max/README.md | 17 -- .../basic/vector_reduce_max/run_vck5000.lit | 11 -- .../basic/vector_reduce_max/test_vck5000.cpp | 148 ------------------ .../basic/vector_reduce_min/Makefile | 21 --- .../basic/vector_reduce_min/README.md | 18 --- .../basic/vector_reduce_min/test_vck5000.cpp | 146 ----------------- .../basic/vector_vector_add/Makefile | 19 +-- .../basic/vector_vector_add/run_vck5000.lit | 9 +- .../basic/vector_vector_mul/Makefile | 21 +-- .../basic/vector_vector_mul/run_vck5000.lit | 9 +- programming_examples/ml/eltwise_add/Makefile | 4 + 20 files changed, 45 insertions(+), 641 deletions(-) delete mode 100644 programming_examples/basic/vector_reduce_add/run_vck5000.lit delete mode 100644 programming_examples/basic/vector_reduce_add/test_vck5000.cpp delete mode 100644 programming_examples/basic/vector_reduce_max/run_vck5000.lit delete mode 100644 programming_examples/basic/vector_reduce_max/test_vck5000.cpp delete mode 100644 programming_examples/basic/vector_reduce_min/test_vck5000.cpp diff --git a/programming_examples/basic/matrix_scalar_add/Makefile b/programming_examples/basic/matrix_scalar_add/Makefile index 82382d45b3..bbeb6a2ed2 100644 --- a/programming_examples/basic/matrix_scalar_add/Makefile +++ b/programming_examples/basic/matrix_scalar_add/Makefile @@ -10,8 +10,6 @@ srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) include ${srcdir}/../../makefile-common -ACDC_AIE = $(dir $(shell which aie-opt))/.. - SHELL := /bin/bash targetname = matrixAddOne @@ -46,17 +44,14 @@ vck5000: col=6 vck5000: build/aie.mlir aiecc.py --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ - -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ - -L/opt/xaiengine/lib \ - -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ - -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ - -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf + -I/opt/xaiengine/include \ + -I${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/include \ + -L/opt/xaiengine/lib \ + -L/lib/x86_64-linux-gnu/ \ + ${srcdir}/test_vck5000.cpp \ + ${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ + -Wl,-R/opt/xaiengine/lib \ + -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf run: ${targetname}.exe build/final.xclbin build/insts.txt ${powershell} ./$< -x build/final.xclbin -i build/insts.txt -k MLIR_AIE diff --git a/programming_examples/basic/matrix_scalar_add/run_vck5000.lit b/programming_examples/basic/matrix_scalar_add/run_vck5000.lit index d314eea2a4..982c28302d 100644 --- a/programming_examples/basic/matrix_scalar_add/run_vck5000.lit +++ b/programming_examples/basic/matrix_scalar_add/run_vck5000.lit @@ -1,9 +1,8 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. +// (c) Copyright 2024 Advanced Micro Devices, Inc. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // // REQUIRES: hsa // -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile vck5000 // RUN: %run_on_vck5000 ./test.elf - diff --git a/programming_examples/basic/passthrough_dmas/Makefile b/programming_examples/basic/passthrough_dmas/Makefile index 04a30c3b16..7b61821cfe 100644 --- a/programming_examples/basic/passthrough_dmas/Makefile +++ b/programming_examples/basic/passthrough_dmas/Makefile @@ -54,15 +54,12 @@ vck5000: col=6 vck5000: build/aie.mlir aiecc.py --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ + -I${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/include \ -L/opt/xaiengine/lib \ -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ + ${srcdir}/test_vck5000.cpp \ + ${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf run_vck5000: diff --git a/programming_examples/basic/passthrough_dmas/run_vck5000.lit b/programming_examples/basic/passthrough_dmas/run_vck5000.lit index 39d4e736a2..982c28302d 100644 --- a/programming_examples/basic/passthrough_dmas/run_vck5000.lit +++ b/programming_examples/basic/passthrough_dmas/run_vck5000.lit @@ -1,8 +1,8 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. +// (c) Copyright 2024 Advanced Micro Devices, Inc. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // // REQUIRES: hsa // -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile vck5000 // RUN: %run_on_vck5000 ./test.elf diff --git a/programming_examples/basic/vector_reduce_add/Makefile b/programming_examples/basic/vector_reduce_add/Makefile index 9f0fabb764..f802b3ae41 100644 --- a/programming_examples/basic/vector_reduce_add/Makefile +++ b/programming_examples/basic/vector_reduce_add/Makefile @@ -10,8 +10,6 @@ srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) include ${srcdir}/../../makefile-common -ACDC_AIE = $(dir $(shell which aie-opt))/.. - targetname = reduce_add devicename = npu col = 0 @@ -54,26 +52,5 @@ trace: clean_trace: rm -rf tmpTrace trace.txt -# Changing variables when we target VCK5000 -vck5000: devicename=xcvc1902 -vck5000: col=6 -vck5000: CHESS_FLAGS=${CHESSCCWRAP1_FLAGS} -vck5000: KERNEL_LIB=${ACDC_AIE}/../aie_kernels/aie1/ - -vck5000: build/aie.mlir build/scale.o - cp build/scale.o* ./ - aiecc.py --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ - -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ - -L/opt/xaiengine/lib \ - -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ - -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ - -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf - clean: clean_trace rm -rf build _build inst aie.mlir.prj core_* test.elf ${targetname}.exe diff --git a/programming_examples/basic/vector_reduce_add/README.md b/programming_examples/basic/vector_reduce_add/README.md index 7548165a1a..238e81496e 100644 --- a/programming_examples/basic/vector_reduce_add/README.md +++ b/programming_examples/basic/vector_reduce_add/README.md @@ -22,8 +22,6 @@ This example does not contain a C++ kernel file. The kernel is expressed in Pyth 1. `test.cpp`: This C++ code is a testbench for the design example targetting Ryzen™ AI (AIE2). The code is responsible for loading the compiled XCLBIN file, configuring the AIE module, providing input data, and executing the AIE design on the NPU. After executing, the program verifies the results. -1. `test_vck5000.cpp`: This C++ code is a testbench for the design example targetting the VCK5000 PCIe card (AIE1). The code is responsible for configuring the AIEs, allocating memory, providing input data, and executing the AIE design on the VCK5000. After executing, the program verifies the results. - ## Ryzen™ AI Usage ### C++ Testbench @@ -41,18 +39,3 @@ To run the design: make run ``` -## VCK5000 Usage - -### C++ Testbench - -To compile the design and C++ testbench: - -``` -make vck5000 -``` - -To run the design: - -``` -./test.elf -``` \ No newline at end of file diff --git a/programming_examples/basic/vector_reduce_add/run_vck5000.lit b/programming_examples/basic/vector_reduce_add/run_vck5000.lit deleted file mode 100644 index 3e162a8ed3..0000000000 --- a/programming_examples/basic/vector_reduce_add/run_vck5000.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: hsa -// -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib -// RUN: %run_on_vck5000 ./test.elf - -// this file needs update -// XFAIL: * diff --git a/programming_examples/basic/vector_reduce_add/test_vck5000.cpp b/programming_examples/basic/vector_reduce_add/test_vck5000.cpp deleted file mode 100644 index 40805520cf..0000000000 --- a/programming_examples/basic/vector_reduce_add/test_vck5000.cpp +++ /dev/null @@ -1,148 +0,0 @@ -//===- test_vck5000.cpp -----------------------------------000---*- C++ -*-===// -// -// This file is licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// Copyright (C) 2024, Advanced Micro Devices, Inc. -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include "memory_allocator.h" -#include "test_library.h" - -#include "aie_data_movement.cpp" -#include "aie_inc.cpp" - -#include "hsa/hsa.h" -#include "hsa/hsa_ext_amd.h" - -constexpr int DMA_COUNT = 64; - -void hsa_check_status(const std::string func_name, hsa_status_t status) { - if (status != HSA_STATUS_SUCCESS) { - const char *status_string(new char[1024]); - hsa_status_string(status, &status_string); - std::cout << func_name << " failed: " << status_string << std::endl; - delete[] status_string; - } else { - std::cout << func_name << " success" << std::endl; - } -} - -int main(int argc, char *argv[]) { - uint64_t row = 0; - uint64_t col = 6; - - std::vector queues; - uint32_t aie_max_queue_size(0); - - aie_libxaie_ctx_t *xaie = mlir_aie_init_libxaie(); - - // This is going to initialize HSA, create a queue - // and get an agent - int ret = mlir_aie_init_device(xaie); - - if (ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" - << std::endl; - return -1; - } - - // Getting access to all of the HSA agents - std::vector agents = xaie->agents; - - if (agents.empty()) { - std::cout << "No agents found. Exiting." << std::endl; - return -1; - } - - std::cout << "Found " << agents.size() << " agents" << std::endl; - - hsa_queue_t *q = xaie->cmd_queue; - - // Adding to our vector of queues - queues.push_back(q); - assert(queues.size() > 0 && "No queues were sucesfully created!"); - - mlir_aie_configure_cores(xaie); - mlir_aie_configure_switchboxes(xaie); - mlir_aie_initialize_locks(xaie); - mlir_aie_configure_dmas(xaie); - mlir_aie_start_cores(xaie); - - // Allocating some device memory - ext_mem_model_t buf0, buf1, buf2; - uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); - uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( - xaie, buf2, 4 /* For some reason can't do 1 */); - mlir_aie_sync_mem_dev(buf0); - mlir_aie_sync_mem_dev(buf1); - mlir_aie_sync_mem_dev(buf2); - - if (in_a == nullptr || in_b == nullptr || out == nullptr) { - std::cout << "Could not allocate in device memory" << std::endl; - return -1; - } - - out[0] = 0xdeface; - for (int i = 0; i < DMA_COUNT; i++) { - in_a[i] = i + 1; - } - - in_a[DMA_COUNT / 2] = 123456; - in_a[DMA_COUNT - 1] = 100; - - // printf("[EDDIE DEBUG] max_val before data movement is %d\n", - // mlir_aie_read_buffer_max_val(xaie, 0)); - - // Pass arguments in the order of dma_memcpys in the mlir - invoke_data_movement(queues[0], &agents[0], out, in_a); - - int errors = 0; - - uint32_t max_val = 0; - for (int i = 0; i < DMA_COUNT; i++) { - uint32_t s = in_a[i]; - if (max_val < s) { - max_val = s; - } - } - - // printf("[EDDIE DEBUG] max_val before data movement is %d\n", - // mlir_aie_read_buffer_max_val(xaie, 0)); - - if (*out != max_val) { - errors++; - printf("[ERROR] Maximum value is %d but kernel returned %d\n", max_val, - *out); - } - - // destroying the queue - hsa_queue_destroy(queues[0]); - - // Shutdown AIR and HSA - mlir_aie_deinit_libxaie(xaie); - - if (!errors) { - printf("PASS!\n"); - return 0; - } else { - printf("fail %d/%d.\n", errors, 1); - return -1; - } -} diff --git a/programming_examples/basic/vector_reduce_max/Makefile b/programming_examples/basic/vector_reduce_max/Makefile index 0c22b7afea..fb0ce195be 100755 --- a/programming_examples/basic/vector_reduce_max/Makefile +++ b/programming_examples/basic/vector_reduce_max/Makefile @@ -54,26 +54,5 @@ trace: clean_trace: rm -rf tmpTrace trace.txt -# Changing variables when we target VCK5000 -vck5000: devicename=xcvc1902 -vck5000: col=6 -vck5000: CHESS_FLAGS=${CHESSCCWRAP1_FLAGS} -vck5000: KERNEL_LIB=${ACDC_AIE}/../aie_kernels/aie1/ - -vck5000: build/aie.mlir build/scale.o - cp build/scale.o* ./ - aiecc.py --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ - -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ - -L/opt/xaiengine/lib \ - -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ - -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ - -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf - clean: clean_trace rm -rf build _build inst aie.mlir.prj core_* test.elf ${targetname}.exe diff --git a/programming_examples/basic/vector_reduce_max/README.md b/programming_examples/basic/vector_reduce_max/README.md index ac2756f2dc..1f935f6823 100644 --- a/programming_examples/basic/vector_reduce_max/README.md +++ b/programming_examples/basic/vector_reduce_max/README.md @@ -22,8 +22,6 @@ This example does not contain a C++ kernel file. The kernel is expressed in Pyth 1. `test.cpp`: This C++ code is a testbench for the design example targetting Ryzen™ AI (AIE2). The code is responsible for loading the compiled XCLBIN file, configuring the AIE module, providing input data, and executing the AIE design on the NPU. After executing, the program verifies the results. -1. `test_vck5000.cpp`: This C++ code is a testbench for the design example targetting the VCK5000 PCIe card (AIE1). The code is responsible for configuring the AIEs, allocating memory, providing input data, and executing the AIE design on the VCK5000. After executing, the program verifies the results. - ## Ryzen™ AI Usage ### C++ Testbench @@ -41,18 +39,3 @@ To run the design: make run ``` -## VCK5000 Usage - -### C++ Testbench - -To compile the design and C++ testbench: - -``` -make vck5000 -``` - -To run the design: - -``` -./test.elf -``` \ No newline at end of file diff --git a/programming_examples/basic/vector_reduce_max/run_vck5000.lit b/programming_examples/basic/vector_reduce_max/run_vck5000.lit deleted file mode 100644 index 3e162a8ed3..0000000000 --- a/programming_examples/basic/vector_reduce_max/run_vck5000.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: hsa -// -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib -// RUN: %run_on_vck5000 ./test.elf - -// this file needs update -// XFAIL: * diff --git a/programming_examples/basic/vector_reduce_max/test_vck5000.cpp b/programming_examples/basic/vector_reduce_max/test_vck5000.cpp deleted file mode 100644 index 40805520cf..0000000000 --- a/programming_examples/basic/vector_reduce_max/test_vck5000.cpp +++ /dev/null @@ -1,148 +0,0 @@ -//===- test_vck5000.cpp -----------------------------------000---*- C++ -*-===// -// -// This file is licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// Copyright (C) 2024, Advanced Micro Devices, Inc. -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include "memory_allocator.h" -#include "test_library.h" - -#include "aie_data_movement.cpp" -#include "aie_inc.cpp" - -#include "hsa/hsa.h" -#include "hsa/hsa_ext_amd.h" - -constexpr int DMA_COUNT = 64; - -void hsa_check_status(const std::string func_name, hsa_status_t status) { - if (status != HSA_STATUS_SUCCESS) { - const char *status_string(new char[1024]); - hsa_status_string(status, &status_string); - std::cout << func_name << " failed: " << status_string << std::endl; - delete[] status_string; - } else { - std::cout << func_name << " success" << std::endl; - } -} - -int main(int argc, char *argv[]) { - uint64_t row = 0; - uint64_t col = 6; - - std::vector queues; - uint32_t aie_max_queue_size(0); - - aie_libxaie_ctx_t *xaie = mlir_aie_init_libxaie(); - - // This is going to initialize HSA, create a queue - // and get an agent - int ret = mlir_aie_init_device(xaie); - - if (ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" - << std::endl; - return -1; - } - - // Getting access to all of the HSA agents - std::vector agents = xaie->agents; - - if (agents.empty()) { - std::cout << "No agents found. Exiting." << std::endl; - return -1; - } - - std::cout << "Found " << agents.size() << " agents" << std::endl; - - hsa_queue_t *q = xaie->cmd_queue; - - // Adding to our vector of queues - queues.push_back(q); - assert(queues.size() > 0 && "No queues were sucesfully created!"); - - mlir_aie_configure_cores(xaie); - mlir_aie_configure_switchboxes(xaie); - mlir_aie_initialize_locks(xaie); - mlir_aie_configure_dmas(xaie); - mlir_aie_start_cores(xaie); - - // Allocating some device memory - ext_mem_model_t buf0, buf1, buf2; - uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); - uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( - xaie, buf2, 4 /* For some reason can't do 1 */); - mlir_aie_sync_mem_dev(buf0); - mlir_aie_sync_mem_dev(buf1); - mlir_aie_sync_mem_dev(buf2); - - if (in_a == nullptr || in_b == nullptr || out == nullptr) { - std::cout << "Could not allocate in device memory" << std::endl; - return -1; - } - - out[0] = 0xdeface; - for (int i = 0; i < DMA_COUNT; i++) { - in_a[i] = i + 1; - } - - in_a[DMA_COUNT / 2] = 123456; - in_a[DMA_COUNT - 1] = 100; - - // printf("[EDDIE DEBUG] max_val before data movement is %d\n", - // mlir_aie_read_buffer_max_val(xaie, 0)); - - // Pass arguments in the order of dma_memcpys in the mlir - invoke_data_movement(queues[0], &agents[0], out, in_a); - - int errors = 0; - - uint32_t max_val = 0; - for (int i = 0; i < DMA_COUNT; i++) { - uint32_t s = in_a[i]; - if (max_val < s) { - max_val = s; - } - } - - // printf("[EDDIE DEBUG] max_val before data movement is %d\n", - // mlir_aie_read_buffer_max_val(xaie, 0)); - - if (*out != max_val) { - errors++; - printf("[ERROR] Maximum value is %d but kernel returned %d\n", max_val, - *out); - } - - // destroying the queue - hsa_queue_destroy(queues[0]); - - // Shutdown AIR and HSA - mlir_aie_deinit_libxaie(xaie); - - if (!errors) { - printf("PASS!\n"); - return 0; - } else { - printf("fail %d/%d.\n", errors, 1); - return -1; - } -} diff --git a/programming_examples/basic/vector_reduce_min/Makefile b/programming_examples/basic/vector_reduce_min/Makefile index 9c6b35a012..6558b10dbc 100755 --- a/programming_examples/basic/vector_reduce_min/Makefile +++ b/programming_examples/basic/vector_reduce_min/Makefile @@ -54,26 +54,5 @@ trace: clean_trace: rm -rf tmpTrace trace.txt -# Changing variables when we target VCK5000 -vck5000: devicename=xcvc1902 -vck5000: col=6 -vck5000: CHESS_FLAGS=${CHESSCCWRAP1_FLAGS} -vck5000: KERNEL_LIB=${ACDC_AIE}/../aie_kernels/aie1/ - -vck5000: build/aie.mlir build/scale.o - cp build/scale.o* ./ - aiecc.py --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ - -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ - -L/opt/xaiengine/lib \ - -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ - -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ - -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf - clean: clean_trace rm -rf build _build inst aie.mlir.prj core_* test.elf ${targetname}.exe diff --git a/programming_examples/basic/vector_reduce_min/README.md b/programming_examples/basic/vector_reduce_min/README.md index feb6328142..a851d97dee 100644 --- a/programming_examples/basic/vector_reduce_min/README.md +++ b/programming_examples/basic/vector_reduce_min/README.md @@ -22,8 +22,6 @@ This example does not contain a C++ kernel file. The kernel is expressed in Pyth 1. `test.cpp`: This C++ code is a testbench for the design example targetting Ryzen™ AI (AIE2). The code is responsible for loading the compiled XCLBIN file, configuring the AIE module, providing input data, and executing the AIE design on the NPU. After executing, the program verifies the results. -1. `test_vck5000.cpp`: This C++ code is a testbench for the design example targetting the VCK5000 PCIe card (AIE1). The code is responsible for configuring the AIEs, allocating memory, providing input data, and executing the AIE design on the VCK5000. After executing, the program verifies the results. - ## Ryzen™ AI Usage ### C++ Testbench @@ -40,19 +38,3 @@ To run the design: ``` make run ``` - -## VCK5000 Usage - -### C++ Testbench - -To compile the design and C++ testbench: - -``` -make vck5000 -``` - -To run the design: - -``` -./test.elf -``` \ No newline at end of file diff --git a/programming_examples/basic/vector_reduce_min/test_vck5000.cpp b/programming_examples/basic/vector_reduce_min/test_vck5000.cpp deleted file mode 100644 index 22ce229f56..0000000000 --- a/programming_examples/basic/vector_reduce_min/test_vck5000.cpp +++ /dev/null @@ -1,146 +0,0 @@ -//===- test.cpp -------------------------------------------------*- C++ -*-===// -// -// Copyright (C) 2020-2022, Xilinx Inc. -// Copyright (C) 2022, Advanced Micro Devices, Inc. -// SPDX-License-Identifier: MIT -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include "memory_allocator.h" -#include "test_library.h" - -#include "aie_data_movement.cpp" -#include "aie_inc.cpp" - -#include "hsa/hsa.h" -#include "hsa/hsa_ext_amd.h" - -constexpr int DMA_COUNT = 64; - -void hsa_check_status(const std::string func_name, hsa_status_t status) { - if (status != HSA_STATUS_SUCCESS) { - const char *status_string(new char[1024]); - hsa_status_string(status, &status_string); - std::cout << func_name << " failed: " << status_string << std::endl; - delete[] status_string; - } else { - std::cout << func_name << " success" << std::endl; - } -} - -int main(int argc, char *argv[]) { - uint64_t row = 0; - uint64_t col = 6; - - std::vector queues; - uint32_t aie_max_queue_size(0); - - aie_libxaie_ctx_t *xaie = mlir_aie_init_libxaie(); - - // This is going to initialize HSA, create a queue - // and get an agent - int ret = mlir_aie_init_device(xaie); - - if (ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" - << std::endl; - return -1; - } - - // Getting access to all of the HSA agents - std::vector agents = xaie->agents; - - if (agents.empty()) { - std::cout << "No agents found. Exiting." << std::endl; - return -1; - } - - std::cout << "Found " << agents.size() << " agents" << std::endl; - - hsa_queue_t *q = xaie->cmd_queue; - - // Adding to our vector of queues - queues.push_back(q); - assert(queues.size() > 0 && "No queues were sucesfully created!"); - - mlir_aie_configure_cores(xaie); - mlir_aie_configure_switchboxes(xaie); - mlir_aie_initialize_locks(xaie); - mlir_aie_configure_dmas(xaie); - mlir_aie_start_cores(xaie); - - // Allocating some device memory - ext_mem_model_t buf0, buf1, buf2; - uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); - uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( - xaie, buf2, 4 /* For some reason can't do 1 */); - mlir_aie_sync_mem_dev(buf0); - mlir_aie_sync_mem_dev(buf1); - mlir_aie_sync_mem_dev(buf2); - - if (in_a == nullptr || in_b == nullptr || out == nullptr) { - std::cout << "Could not allocate in device memory" << std::endl; - return -1; - } - - out[0] = 0xdeface; - for (int i = 0; i < DMA_COUNT; i++) { - in_a[i] = i + 1; - } - - in_a[DMA_COUNT / 2] = 123456; - in_a[DMA_COUNT - 1] = 100; - - // printf("[EDDIE DEBUG] min_val before data movement is %d\n", - // mlir_aie_read_buffer_min_val(xaie, 0)); - - // Pass arguments in the order of dma_memcpys in the mlir - invoke_data_movement(queues[0], &agents[0], out, in_a); - - int errors = 0; - - uint32_t min_val = INT_MAX; - for (int i = 0; i < DMA_COUNT; i++) { - uint32_t s = in_a[i]; - if (s < min_val) { - min_val = s; - } - } - - // printf("[EDDIE DEBUG] min_val before data movement is %d\n", - // mlir_aie_read_buffer_min_val(xaie, 0)); - - if (*out != min_val) { - errors++; - printf("[ERROR] Minimum value is %d but kernel returned %d\n", min_val, - *out); - } - - // destroying the queue - hsa_queue_destroy(queues[0]); - - // Shutdown AIR and HSA - mlir_aie_deinit_libxaie(xaie); - - if (!errors) { - printf("PASS!\n"); - return 0; - } else { - printf("fail %d/%d.\n", errors, 1); - return -1; - } -} diff --git a/programming_examples/basic/vector_vector_add/Makefile b/programming_examples/basic/vector_vector_add/Makefile index 9d732bb0e3..8687ff97e2 100755 --- a/programming_examples/basic/vector_vector_add/Makefile +++ b/programming_examples/basic/vector_vector_add/Makefile @@ -44,17 +44,14 @@ vck5000: col=6 vck5000: build/aie.mlir aiecc.py --xchesscc --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ - -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ - -L/opt/xaiengine/lib \ - -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ - -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ - -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf + -I/opt/xaiengine/include \ + -I${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/include \ + -L/opt/xaiengine/lib \ + -L/lib/x86_64-linux-gnu/ \ + ${srcdir}/test_vck5000.cpp \ + ${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ + -Wl,-R/opt/xaiengine/lib \ + -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf run: ${targetname}.exe build/final.xclbin build/insts.txt ${powershell} ./$< -x build/final.xclbin -i build/insts.txt -k MLIR_AIE diff --git a/programming_examples/basic/vector_vector_add/run_vck5000.lit b/programming_examples/basic/vector_vector_add/run_vck5000.lit index 02055c5eca..dcaa9f99c5 100644 --- a/programming_examples/basic/vector_vector_add/run_vck5000.lit +++ b/programming_examples/basic/vector_vector_add/run_vck5000.lit @@ -1,9 +1,8 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. +// (c) Copyright 2024 Advanced Micro Devices, Inc. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // // REQUIRES: hsa, chess // -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py --xchesscc %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib -// RUN: %run_on_vck5000 ./test.elf - +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile vck5000 +// RUN: %run_on_vck5000 ./test.elf diff --git a/programming_examples/basic/vector_vector_mul/Makefile b/programming_examples/basic/vector_vector_mul/Makefile index 6697a884ca..35a8054ed1 100755 --- a/programming_examples/basic/vector_vector_mul/Makefile +++ b/programming_examples/basic/vector_vector_mul/Makefile @@ -39,24 +39,19 @@ endif # Changing variables when we target VCK5000 -ACDC_AIE = $(dir $(shell which aie-opt))/.. - vck5000: devicename=xcvc1902 vck5000: col=6 vck5000: build/aie.mlir aiecc.py --xchesscc --link_against_hsa --host-target=x86_64-amd-linux-gnu build/aie.mlir \ - -I/opt/xaiengine/include \ - -I$(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/include \ - -I${ROCM_ROOT}/include \ - -L/opt/xaiengine/lib \ - -L/lib/x86_64-linux-gnu/ \ - test_vck5000.cpp \ - $(ACDC_AIE)/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ - ${ROCM_ROOT}/lib/libhsa-runtime64.so.1.9.0 \ - -Wl,-R/opt/xaiengine/lib \ - -Wl,-rpath,${ROCM_ROOT}/lib \ - -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf + -I/opt/xaiengine/include \ + -I${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/include \ + -L/opt/xaiengine/lib \ + -L/lib/x86_64-linux-gnu/ \ + ${srcdir}/test_vck5000.cpp \ + ${srcdir}/../../../install/runtime_lib/x86_64-hsa/test_lib/src/test_library.cpp \ + -Wl,-R/opt/xaiengine/lib \ + -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf run: ${targetname}.exe build/final.xclbin build/insts.txt diff --git a/programming_examples/basic/vector_vector_mul/run_vck5000.lit b/programming_examples/basic/vector_vector_mul/run_vck5000.lit index 02055c5eca..dcaa9f99c5 100644 --- a/programming_examples/basic/vector_vector_mul/run_vck5000.lit +++ b/programming_examples/basic/vector_vector_mul/run_vck5000.lit @@ -1,9 +1,8 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. +// (c) Copyright 2024 Advanced Micro Devices, Inc. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // // REQUIRES: hsa, chess // -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py --xchesscc %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib -// RUN: %run_on_vck5000 ./test.elf - +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile vck5000 +// RUN: %run_on_vck5000 ./test.elf diff --git a/programming_examples/ml/eltwise_add/Makefile b/programming_examples/ml/eltwise_add/Makefile index 3597fc276b..bab7f2c93d 100644 --- a/programming_examples/ml/eltwise_add/Makefile +++ b/programming_examples/ml/eltwise_add/Makefile @@ -53,6 +53,10 @@ else cp _build/${targetname} $@ endif +build/aie.mlir: aie2.py + mkdir -p ${@D} + python3 $< ${devicename} ${col} > $@ + run: ${targetname}.exe build/final.xclbin build/insts.txt ${powershell} ./$< -x build/final.xclbin -i build/insts.txt -k MLIR_AIE