diff --git a/programming_examples/basic/dma_transpose/Makefile b/programming_examples/basic/dma_transpose/Makefile index 39274a8a48..1cc5f1e83a 100644 --- a/programming_examples/basic/dma_transpose/Makefile +++ b/programming_examples/basic/dma_transpose/Makefile @@ -6,7 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common SHELL := /bin/bash @@ -16,27 +18,19 @@ targetname = dmaTranspose M ?= 64 K ?= 32 -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< ${M} ${K} > $@ -.PHONY: inst/insts.txt -inst/insts.txt: aie2.py - rm -rf inst - mkdir -p inst - python3 $< ${LENGTH} > inst/aie.mlir - pushd inst && aiecc.py --aie-only-generate-npu --npu-insts-name=insts.txt aie.mlir && popd - ${powershell} ./build/${targetname}.exe -x build/final.xclbin -i inst/insts.txt -k MLIR_AIE -l ${LENGTH} - build/final.xclbin: build/aie.mlir mkdir -p ${@D} cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts.txt $(<:%=../%) -${targetname}.exe: test.cpp +${targetname}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ diff --git a/programming_examples/basic/dma_transpose/run.lit b/programming_examples/basic/dma_transpose/run.lit deleted file mode 100644 index d513e9bca8..0000000000 --- a/programming_examples/basic/dma_transpose/run.lit +++ /dev/null @@ -1,10 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: %python %S/aie2.py 64 32 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt --M 64 --K 32 | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/dma_transpose/run_makefile.lit b/programming_examples/basic/dma_transpose/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/dma_transpose/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/matrix_add_one/run.lit b/programming_examples/basic/matrix_add_one/run.lit deleted file mode 100644 index 1922c01828..0000000000 --- a/programming_examples/basic/matrix_add_one/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai -// -// RUN: %python %S/aie2.py npu 0 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! - diff --git a/programming_examples/basic/matrix_multiplication/makefile-common b/programming_examples/basic/matrix_multiplication/makefile-common index c8486817a0..7f361e3022 100644 --- a/programming_examples/basic/matrix_multiplication/makefile-common +++ b/programming_examples/basic/matrix_multiplication/makefile-common @@ -27,11 +27,12 @@ # - M, K, N -- (optional) dimensions of matrices, may be used by design; # N=1 for matrix-vector +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) #include ${CURDIR}/../../makefile-common current_dir := $(dir $(abspath $(lastword $(MAKEFILE_LIST)))) #include ${current_dir}../../makefile-common SELF_DIR := $(dir $(lastword $(MAKEFILE_LIST))) -include ${SELF_DIR}../../makefile-common +include ${current_dir}../../makefile-common # defaults; overwrite if needed M?=512 @@ -46,16 +47,16 @@ insts_target?=build/insts_${M}x${K}x${N}.txt runargs?=-v 1 --warmup 10 --iters 10 -kernels_dir=../../../../aie_kernels/aie2 +kernels_dir=${srcdir}/../../../../aie_kernels/aie2 .PHONY: all all: ${xclbin_target} ${insts_target} ${targetname}.exe build/%.o: ${kernels_dir}/%.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -DBIT_WIDTH=8 -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -DBIT_WIDTH=8 -c $< -o ${@F} -${mlir_target}: aie2.py +${mlir_target}: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< -M $M -K $K -N $N > $@ @@ -64,10 +65,10 @@ ${xclbin_target}: ${mlir_target} ${kernels:%=build/%.o} cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=${insts_target:build/%=%} $(<:%=../%) -${targetname}.exe: test.cpp ../test.cpp ../common.h +${targetname}.exe: ${srcdir}/test.cpp ${srcdir}/../test.cpp ${srcdir}/../common.h rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake -E env CXXFLAGS="-std=c++23 -ggdb" cmake ../.. -D CMAKE_C_COMPILER=gcc-13 -D CMAKE_CXX_COMPILER=g++-13 -DTARGET_NAME=${targetname} -Dsubdir=${subdir} + cd _build && ${powershell} cmake -E env CXXFLAGS="-std=c++23 -ggdb" cmake ${srcdir}/.. -D CMAKE_C_COMPILER=gcc-13 -D CMAKE_CXX_COMPILER=g++-13 -DTARGET_NAME=${targetname} -Dsubdir=${subdir} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ diff --git a/programming_examples/basic/matrix_multiplication/matrix_vector/run.lit b/programming_examples/basic/matrix_multiplication/matrix_vector/run.lit deleted file mode 100644 index eeaa69352a..0000000000 --- a/programming_examples/basic/matrix_multiplication/matrix_vector/run.lit +++ /dev/null @@ -1,12 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../../aie_kernels/aie2/mv.cc -o ./mv.o -// RUN: %python %S/aie2.py -M 288 -K 288 -N 1 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++-13 %S/test.cpp -o test.exe -std=c++23 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt -M 288 -K 288 -N 1 -v 1 | FileCheck %s -// CHECK: PASS! - diff --git a/programming_examples/basic/matrix_multiplication/matrix_vector/run_makefile.lit b/programming_examples/basic/matrix_multiplication/matrix_vector/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/matrix_multiplication/matrix_vector/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/matrix_multiplication/single_core/run.lit b/programming_examples/basic/matrix_multiplication/single_core/run.lit deleted file mode 100644 index 6f6a32320a..0000000000 --- a/programming_examples/basic/matrix_multiplication/single_core/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../../aie_kernels/aie2/mm.cc -o ./mm.o -// RUN: %python %S/aie2.py -M 256 -K 256 -N 256 > ./aie.mlir -// RUN: %python aiecc.py --xbridge --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++-13 %S/test.cpp -o test.exe -std=c++23 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt -M 256 -K 256 -N 256 -v 1 | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/matrix_multiplication/single_core/run_makefile.lit b/programming_examples/basic/matrix_multiplication/single_core/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/matrix_multiplication/single_core/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/matrix_multiplication/whole_array/Makefile b/programming_examples/basic/matrix_multiplication/whole_array/Makefile index f4a06812ca..2289d762c6 100644 --- a/programming_examples/basic/matrix_multiplication/whole_array/Makefile +++ b/programming_examples/basic/matrix_multiplication/whole_array/Makefile @@ -5,6 +5,7 @@ # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception # ##===----------------------------------------------------------------------===## +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) subdir=whole_array targetname=matrixMultiplication @@ -14,4 +15,4 @@ M?=512 K?=512 N?=512 -include ../makefile-common +include ${srcdir}/../makefile-common diff --git a/programming_examples/basic/matrix_multiplication/whole_array/run.lit b/programming_examples/basic/matrix_multiplication/whole_array/run.lit deleted file mode 100644 index 0a15bd164e..0000000000 --- a/programming_examples/basic/matrix_multiplication/whole_array/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../../aie_kernels/aie2/mm.cc -o ./mm.o -// RUN: %python %S/aie2.py -M 512 -K 512 -N 512 > ./aie.mlir -// RUN: %python aiecc.py --xbridge --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++-13 %S/test.cpp -o test.exe -std=c++23 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt -v 1 -M 512 -K 512 -N 512 | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/matrix_multiplication/whole_array/run_makefile.lit b/programming_examples/basic/matrix_multiplication/whole_array/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/matrix_multiplication/whole_array/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_mult/CMakeLists.txt b/programming_examples/basic/matrix_scalar_add/CMakeLists.txt similarity index 97% rename from programming_examples/basic/vector_mult/CMakeLists.txt rename to programming_examples/basic/matrix_scalar_add/CMakeLists.txt index ee2050a94e..20f5d8a4a3 100644 --- a/programming_examples/basic/vector_mult/CMakeLists.txt +++ b/programming_examples/basic/matrix_scalar_add/CMakeLists.txt @@ -49,8 +49,6 @@ add_executable(${currentTarget} test.cpp ) -set_property(TARGET ${currentTarget} PROPERTY CXX_STANDARD 23) - target_compile_definitions(${currentTarget} PUBLIC DISABLE_ABI_CHECK=1) target_include_directories (${currentTarget} PUBLIC diff --git a/programming_examples/basic/matrix_add_one/Makefile b/programming_examples/basic/matrix_scalar_add/Makefile similarity index 88% rename from programming_examples/basic/matrix_add_one/Makefile rename to programming_examples/basic/matrix_scalar_add/Makefile index 83014fbeaf..82382d45b3 100644 --- a/programming_examples/basic/matrix_add_one/Makefile +++ b/programming_examples/basic/matrix_scalar_add/Makefile @@ -6,7 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common ACDC_AIE = $(dir $(shell which aie-opt))/.. @@ -23,10 +25,10 @@ build/final.xclbin: build/aie.mlir cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ --xclbin-name=${@F} --npu-insts-name=insts.txt ${ $@ diff --git a/programming_examples/basic/matrix_add_one/README.md b/programming_examples/basic/matrix_scalar_add/README.md similarity index 95% rename from programming_examples/basic/matrix_add_one/README.md rename to programming_examples/basic/matrix_scalar_add/README.md index 8516bdfe47..1b50d4af86 100644 --- a/programming_examples/basic/matrix_add_one/README.md +++ b/programming_examples/basic/matrix_scalar_add/README.md @@ -8,7 +8,7 @@ // //===----------------------------------------------------------------------===//--> -# Matrix Addition +# Matrix Scalar Addition Single tile performs a very simple `+` operation where the kernel loads data from local memory, increments the value by `1` and stores it back. The DMA in the Shim tile is programmed to bring the bottom left `8x16` portion of a larger `16x128` matrix into the tile to perform the operation. This reference design can be run on either a RyzenAI NPU or a VCK5000. diff --git a/programming_examples/basic/matrix_add_one/aie2.py b/programming_examples/basic/matrix_scalar_add/aie2.py similarity index 100% rename from programming_examples/basic/matrix_add_one/aie2.py rename to programming_examples/basic/matrix_scalar_add/aie2.py diff --git a/programming_examples/basic/matrix_scalar_add/run_makefile.lit b/programming_examples/basic/matrix_scalar_add/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/matrix_scalar_add/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/matrix_add_one/run_vck5000.lit b/programming_examples/basic/matrix_scalar_add/run_vck5000.lit similarity index 100% rename from programming_examples/basic/matrix_add_one/run_vck5000.lit rename to programming_examples/basic/matrix_scalar_add/run_vck5000.lit diff --git a/programming_examples/basic/matrix_add_one/test.cpp b/programming_examples/basic/matrix_scalar_add/test.cpp similarity index 100% rename from programming_examples/basic/matrix_add_one/test.cpp rename to programming_examples/basic/matrix_scalar_add/test.cpp diff --git a/programming_examples/basic/matrix_add_one/test_vck5000.cpp b/programming_examples/basic/matrix_scalar_add/test_vck5000.cpp similarity index 100% rename from programming_examples/basic/matrix_add_one/test_vck5000.cpp rename to programming_examples/basic/matrix_scalar_add/test_vck5000.cpp diff --git a/programming_examples/basic/passthrough_dmas/Makefile b/programming_examples/basic/passthrough_dmas/Makefile index e09c8a91aa..04a30c3b16 100644 --- a/programming_examples/basic/passthrough_dmas/Makefile +++ b/programming_examples/basic/passthrough_dmas/Makefile @@ -6,26 +6,25 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) -ACDC_AIE = $(dir $(shell which aie-opt))/.. - -SHELL := /bin/bash +include ${srcdir}/../../makefile-common all: build/final.xclbin build/insts.txt +devicename ?= npu targetname = passThroughDMAs LENGTH ?= 4096 -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} - python3 $< ${devicename} ${col} ${LENGTH} ${LENGTH} > $@ + python3 $< ${LENGTH} ${devicename} ${col} > $@ .PHONY: inst/insts.txt -inst/insts.txt: aie2.py +inst/insts.txt: ${srcdir}/aie2.py rm -rf inst mkdir -p inst - python3 $< ${devicename} ${col} ${LENGTH} > inst/aie.mlir + python3 $< ${LENGTH} > inst/aie.mlir pushd inst && aiecc.py --aie-only-generate-npu --npu-insts-name=insts.txt aie.mlir && popd ${powershell} ./build/${targetname}.exe -x build/final.xclbin -i inst/insts.txt -k MLIR_AIE -l ${LENGTH} @@ -34,10 +33,10 @@ build/final.xclbin: build/aie.mlir cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts.txt $(<:%=../%) -${targetname}.exe: test.cpp +${targetname}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ @@ -66,7 +65,8 @@ vck5000: build/aie.mlir -Wl,-rpath,${ROCM_ROOT}/lib \ -Wl,--whole-archive -Wl,--no-whole-archive -lstdc++ -ldl -lelf -o test.elf - +run_vck5000: + test.elf clean: rm -rf build _build inst ${targetname}.exe diff --git a/programming_examples/basic/passthrough_dmas/aie2.py b/programming_examples/basic/passthrough_dmas/aie2.py index f8dc35a6d9..c07db59a7f 100755 --- a/programming_examples/basic/passthrough_dmas/aie2.py +++ b/programming_examples/basic/passthrough_dmas/aie2.py @@ -14,23 +14,22 @@ from aie.extras.context import mlir_mod_ctx N = 4096 +dev = AIEDevice.npu +col = 0 - -# Deciphering the command line arguments -if len(sys.argv) < 3: - raise ValueError("[ERROR] Need 2 command line arguments (Device name, Col)") - -if len(sys.argv) == 4: +if len(sys.argv) > 1: N = int(sys.argv[1]) -if sys.argv[1] == "npu": - dev = AIEDevice.npu -elif sys.argv[1] == "xcvc1902": - dev = AIEDevice.xcvc1902 -else: - raise ValueError("[ERROR] Device name {} is unknown".format(sys.argv[1])) +if len(sys.argv) > 2: + if sys.argv[2] == "npu": + dev = AIEDevice.npu + elif sys.argv[2] == "xcvc1902": + dev = AIEDevice.xcvc1902 + else: + raise ValueError("[ERROR] Device name {} is unknown".format(sys.argv[2])) -col = int(sys.argv[2]) +if len(sys.argv) > 3: + col = int(sys.argv[3]) def my_passthrough(): diff --git a/programming_examples/basic/passthrough_dmas/run.lit b/programming_examples/basic/passthrough_dmas/run.lit deleted file mode 100644 index a466533551..0000000000 --- a/programming_examples/basic/passthrough_dmas/run.lit +++ /dev/null @@ -1,10 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai -// -// RUN: %python %S/aie2.py npu 0 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt -l 4096 | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/passthrough_dmas/run_makefile.lit b/programming_examples/basic/passthrough_dmas/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/passthrough_dmas/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/passthrough_kernel/Makefile b/programming_examples/basic/passthrough_kernel/Makefile index 458b992521..8ad3b4fe2f 100644 --- a/programming_examples/basic/passthrough_kernel/Makefile +++ b/programming_examples/basic/passthrough_kernel/Makefile @@ -6,35 +6,46 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) -VPATH := ../../../aie_kernels/generic - -PASSTHROUGH_SIZE = 4096 +include ${srcdir}/../../makefile-common targetname = passThroughKernel +VPATH := ${srcdir}/../../../aie_kernels/generic +data_size = 4096 +trace_size = 8192 +PASSTHROUGH_SIZE = ${data_size} .PHONY: all template clean -all: build/final_${PASSTHROUGH_SIZE}.xclbin +all: build/final_${data_size}.xclbin + +build/aie2_lineBased_8b_${data_size}.mlir: ${srcdir}/aie2.py + mkdir -p ${@D} + python3 $< ${data_size} 0 > $@ -build/aie2_lineBased_8b_${PASSTHROUGH_SIZE}.mlir: aie2.py +build/aie_trace__lineBased_8b_${data_size}.mlir: ${srcdir}/aie2.py mkdir -p ${@D} - python3 $< ${PASSTHROUGH_SIZE} > $@ + python3 $< ${data_size} ${trace_size} > $@ build/passThrough.cc.o: passThrough.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -DBIT_WIDTH=8 -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -DBIT_WIDTH=8 -c $< -o ${@F} -build/final_${PASSTHROUGH_SIZE}.xclbin: build/aie2_lineBased_8b_${PASSTHROUGH_SIZE}.mlir build/passThrough.cc.o +build/final_${data_size}.xclbin: build/aie2_lineBased_8b_${data_size}.mlir build/passThrough.cc.o mkdir -p ${@D} cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ - --xclbin-name=${@F} --npu-insts-name=insts.txt $(<:%=../%) + --xclbin-name=${@F} --npu-insts-name=insts_${data_size}.txt $(<:%=../%) -${targetname}.exe: test.cpp +build/final_trace_${data_size}.xclbin: build/aie2_lineBased_8b_${data_size}.mlir build/passThrough.cc.o + mkdir -p ${@D} + cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ + --xclbin-name=${@F} --npu-insts-name=insts_${data_size}.txt $(<:%=../%) + +${targetname}_${data_size}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} -DPASSTHROUGH_SIZE=${PASSTHROUGH_SIZE} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} -DPASSTHROUGH_SIZE=${data_size} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ @@ -42,20 +53,25 @@ else cp _build/${targetname} $@ endif -run: ${targetname}.exe build/final_${PASSTHROUGH_SIZE}.xclbin build/insts.txt - ${powershell} ./$< -x build/final_${PASSTHROUGH_SIZE}.xclbin -i build/insts.txt -k MLIR_AIE +run: ${targetname}_${data_size}.exe build/final_${data_size}.xclbin build/insts_${data_size}.txt + ${powershell} ./$< -x build/final_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE + +#run-g: ${targetname}.exe build/final_${data_size}.xclbin build/insts.txt +# ${powershell} ./$< -x build/final_${data_size}.xclbin -i build/insts.txt -k MLIR_AIE -t ${trace_size} -run-g: ${targetname}.exe build/final_${PASSTHROUGH_SIZE}.xclbin build/insts.txt - ${powershell} ./$< -x build/final_${PASSTHROUGH_SIZE}.xclbin -i build/insts.txt -k MLIR_AIE -t 8192 +run_py: build/final_${data_size}.xclbin build/insts_${data_size}.txt + ${powershell} python3 ${srcdir}/test.py -s ${data_size} -x build/final_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -run_py: build/final_${PASSTHROUGH_SIZE}.xclbin build/insts.txt - ${powershell} python3 test.py -s ${PASSTHROUGH_SIZE} -x build/final_${PASSTHROUGH_SIZE}.xclbin -i build/insts.txt -k MLIR_AIE +trace: ${targetname}_${data_size}.exe build/final_trace_${data_size}.xclbin build/insts_${data_size}.txt + ${powershell} ./$< -x build/final_trace_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -t ${trace_size} + ../../utils/parse_trace.py --filename trace.txt --mlir build/aie_trace_${data_size}.mlir --colshift 1 > trace_vs.json -trace: - ../../utils/parse_eventIR.py --filename trace.txt --mlir build/aie.mlir --colshift 1 > parse_eventIR_vs.json +trace_py: build/final_trace_${data_size}.xclbin build/insts_${data_size}.txt + ${powershell} python3 ${srcdir}/test.py -x build/final_trace_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -t ${trace_size} -s ${data_size} + ../../utils/parse_trace.py --filename trace.txt --mlir build/aie_trace_${data_size}.mlir --colshift 1 > trace_vs.json clean_trace: - rm -rf tmpTrace trace.txt + rm -rf tmpTrace trace.txt parse*json trace*json clean: - rm -rf build _build ${targetname}.exe + rm -rf build _build ${targetname}*.exe diff --git a/programming_examples/basic/passthrough_kernel/aie2.py b/programming_examples/basic/passthrough_kernel/aie2.py index ddde187360..a9268dc845 100644 --- a/programming_examples/basic/passthrough_kernel/aie2.py +++ b/programming_examples/basic/passthrough_kernel/aie2.py @@ -14,20 +14,11 @@ import aie.utils.trace as trace_utils -N = 1024 -if len(sys.argv) == 2: - N = int(sys.argv[1]) - -lineWidthInBytes = N // 4 # chop input in 4 sub-tensors -lineWidthInInt32s = lineWidthInBytes // 4 - -enableTrace = False -traceSizeInBytes = 8192 -traceSizeInInt32s = traceSizeInBytes // 4 - - -def passthroughKernel(): +def passthroughKernel(vector_size, trace_size): + N = vector_size + lineWidthInBytes = N // 4 # chop input in 4 sub-tensors + lineWidthInInt32s = lineWidthInBytes // 4 @device(AIEDevice.npu) def device_body(): @@ -43,7 +34,8 @@ def device_body(): ShimTile = tile(0, 0) ComputeTile2 = tile(0, 2) - if enableTrace: + # Set up a circuit-switched flow from core to shim for tracing information + if trace_size > 0: flow(ComputeTile2, WireBundle.Trace, 0, ShimTile, WireBundle.DMA, 1) # AIE-array data movement with object fifos @@ -67,22 +59,17 @@ def core_body(): tensorSize = N tensorSizeInInt32s = tensorSize // 4 - tensor_ty = T.memref(lineWidthInInt32s, T.i32()) + tensor_ty = T.memref(tensorSizeInInt32s, T.i32()) @FuncOp.from_py_func(tensor_ty, tensor_ty, tensor_ty) def sequence(inTensor, outTensor, notUsed): - if enableTrace: + if trace_size > 0: trace_utils.configure_simple_tracing_aie2( ComputeTile2, ShimTile, - channel=1, - bd_id=13, ddr_id=1, - size=traceSizeInBytes, + size=trace_size, offset=tensorSize, - start=0x1, - stop=0x0, - events=[0x4B, 0x22, 0x21, 0x25, 0x2D, 0x2C, 0x1A, 0x4F], ) npu_dma_memcpy_nd( @@ -100,6 +87,14 @@ def sequence(inTensor, outTensor, notUsed): npu_sync(column=0, row=0, direction=0, channel=0) +try: + vector_size = int(sys.argv[1]) + if vector_size % 64 != 0 or vector_size < 512: + print("Vector size must be a multiple of 64 and greater than or equal to 512") + raise ValueError + trace_size = 0 if (len(sys.argv) != 3) else int(sys.argv[2]) +except ValueError: + print("Argument has inappropriate value") with mlir_mod_ctx() as ctx: - passthroughKernel() + passthroughKernel(vector_size, trace_size) print(ctx.module) diff --git a/programming_examples/basic/passthrough_kernel/run.lit b/programming_examples/basic/passthrough_kernel/run.lit deleted file mode 100644 index 7f1c2318b2..0000000000 --- a/programming_examples/basic/passthrough_kernel/run.lit +++ /dev/null @@ -1,12 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -DBIT_WIDTH=8 -c %S/../../../aie_kernels/generic/passThrough.cc -o passThrough.cc.o -// RUN: %python %S/aie2.py 4096 | aie-opt -cse -canonicalize -o ./aie.mlir -// RUN: %python aiecc.py --xbridge --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++ %S/test.cpp -o test.exe -std=c++23 -Wall -DPASSTHROUGH_SIZE=4096 -I%S/../../../runtime_lib/test_lib %S/../../../runtime_lib/test_lib/test_utils.cpp %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// RUN: %run_on_npu %python %S/test.py -x aie.xclbin -i insts.txt -k MLIR_AIE -s 4096 | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/passthrough_kernel/run_makefile.lit b/programming_examples/basic/passthrough_kernel/run_makefile.lit new file mode 100644 index 0000000000..eb23a0be03 --- /dev/null +++ b/programming_examples/basic/passthrough_kernel/run_makefile.lit @@ -0,0 +1,10 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// RUN: %run_on_npu make -f %S/Makefile run_py | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_add/run.lit b/programming_examples/basic/vector_add/run.lit deleted file mode 100644 index 1922c01828..0000000000 --- a/programming_examples/basic/vector_add/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai -// -// RUN: %python %S/aie2.py npu 0 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! - diff --git a/programming_examples/basic/vector_exp/Makefile b/programming_examples/basic/vector_exp/Makefile index 751779d2ee..82325049ae 100644 --- a/programming_examples/basic/vector_exp/Makefile +++ b/programming_examples/basic/vector_exp/Makefile @@ -6,26 +6,28 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common targetname = testExp all: build/final.xclbin build/insts.txt -VPATH := ../../../aie_kernels/aie2 +VPATH := ${srcdir}/../../../aie_kernels/aie2 build/exp.o: bf16_exp.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -I../../../../aie_runtime_lib/AIE2 -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -I../../../../aie_runtime_lib/AIE2 -c $< -o ${@F} -build/lut_based_ops.o: ../../../aie_runtime_lib/AIE2/lut_based_ops.cpp +build/lut_based_ops.o: ${srcdir}/../../../aie_runtime_lib/AIE2/lut_based_ops.cpp mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -I. -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -I. -c $< -o ${@F} build/kernels.a: build/exp.o build/lut_based_ops.o ar rvs $@ $+ -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< > $@ @@ -34,10 +36,10 @@ build/final.xclbin: build/aie.mlir build/kernels.a cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts.txt $(<:%=../%) -${targetname}.exe: test.cpp +${targetname}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ diff --git a/programming_examples/basic/vector_exp/run.lit b/programming_examples/basic/vector_exp/run.lit deleted file mode 100644 index 247ca37a33..0000000000 --- a/programming_examples/basic/vector_exp/run.lit +++ /dev/null @@ -1,13 +0,0 @@ -// (c) Copyright 2024 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -I %S/../../../aie_runtime_lib/AIE2 -c %S/../../../aie_kernels/aie2/bf16_exp.cc -o exp.o -// RUN: xchesscc_wrapper aie2 -I %aietools/include -I. -c %S/../../../aie_runtime_lib/AIE2/lut_based_ops.cpp -o lut_based_ops.o -// RUN: ar rvs kernels.a exp.o lut_based_ops.o -// RUN: %python %S/aie2.py npu 0 | aie-opt -cse -canonicalize -o ./aie.mlir -// RUN: %python aiecc.py --xbridge --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++-13 %S/test.cpp -o test.exe -std=c++23 -Wall -I%S/../../../runtime_lib/test_lib %S/../../../runtime_lib/test_lib/test_utils.cpp %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/vector_exp/run_makefile.lit b/programming_examples/basic/vector_exp/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_exp/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_mult/run.lit b/programming_examples/basic/vector_mult/run.lit deleted file mode 100644 index 1922c01828..0000000000 --- a/programming_examples/basic/vector_mult/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai -// -// RUN: %python %S/aie2.py npu 0 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! - diff --git a/programming_examples/basic/vector_reduce_add/Makefile b/programming_examples/basic/vector_reduce_add/Makefile index b9f45c87a7..9f0fabb764 100644 --- a/programming_examples/basic/vector_reduce_add/Makefile +++ b/programming_examples/basic/vector_reduce_add/Makefile @@ -6,7 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common ACDC_AIE = $(dir $(shell which aie-opt))/.. @@ -17,13 +19,13 @@ CHESS_FLAGS=${CHESSCCWRAP2_FLAGS} all: build/final.xclbin build/insts.txt -VPATH := ../../../aie_kernels/aie2 +VPATH := ${srcdir}/../../../aie_kernels/aie2 build/%.cc.o: %.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $< -o ${@F} -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< ${devicename} ${col} > $@ @@ -32,10 +34,10 @@ build/final.xclbin: build/aie.mlir build/reduce_add.cc.o cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts.txt $(<:%=../%) -${targetname}.exe: test.cpp +${targetname}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ diff --git a/programming_examples/basic/vector_reduce_add/run.lit b/programming_examples/basic/vector_reduce_add/run.lit deleted file mode 100644 index f35b24884f..0000000000 --- a/programming_examples/basic/vector_reduce_add/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../aie_kernels/aie2/reduce_add.cc -o reduce_add.cc.o -// RUN: %python %S/aie2.py npu 0 | aie-opt -cse -canonicalize -o ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall -I%S/../../../runtime_lib/test_lib %S/../../../runtime_lib/test_lib/test_utils.cpp %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/vector_reduce_add/run_makefile.lit b/programming_examples/basic/vector_reduce_add/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_reduce_add/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_reduce_max/Makefile b/programming_examples/basic/vector_reduce_max/Makefile index 37bcaf43e8..0c22b7afea 100755 --- a/programming_examples/basic/vector_reduce_max/Makefile +++ b/programming_examples/basic/vector_reduce_max/Makefile @@ -6,7 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common ACDC_AIE = $(dir $(shell which aie-opt))/.. @@ -17,13 +19,13 @@ CHESS_FLAGS=${CHESSCCWRAP2_FLAGS} all: build/final.xclbin build/insts.txt -VPATH := ../../../aie_kernels/aie2 +VPATH := ${srcdir}/../../../aie_kernels/aie2 build/%.cc.o: %.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $< -o ${@F} -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< ${devicename} ${col} > $@ @@ -32,10 +34,10 @@ build/final.xclbin: build/aie.mlir build/reduce_max.cc.o cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts.txt $(<:%=../%) -${targetname}.exe: test.cpp +${targetname}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ diff --git a/programming_examples/basic/vector_reduce_max/run.lit b/programming_examples/basic/vector_reduce_max/run.lit deleted file mode 100644 index 584d7c1628..0000000000 --- a/programming_examples/basic/vector_reduce_max/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../aie_kernels/aie2/reduce_max.cc -o reduce_max.cc.o -// RUN: %python %S/aie2.py npu 0 | aie-opt -cse -canonicalize -o ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++ %S/test.cpp -o test.exe -std=c++23 -Wall -I%S/../../../runtime_lib/test_lib %S/../../../runtime_lib/test_lib/test_utils.cpp %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/vector_reduce_max/run_makefile.lit b/programming_examples/basic/vector_reduce_max/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_reduce_max/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_reduce_min/Makefile b/programming_examples/basic/vector_reduce_min/Makefile index 3f4d1105df..9c6b35a012 100755 --- a/programming_examples/basic/vector_reduce_min/Makefile +++ b/programming_examples/basic/vector_reduce_min/Makefile @@ -6,7 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common ACDC_AIE = $(dir $(shell which aie-opt))/.. @@ -17,13 +19,13 @@ CHESS_FLAGS=${CHESSCCWRAP2_FLAGS} all: build/final.xclbin build/insts.txt -VPATH := ../../../aie_kernels/aie2 +VPATH := ${srcdir}/../../../aie_kernels/aie2 build/%.cc.o: %.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $< -o ${@F} -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< ${devicename} ${col} > $@ @@ -32,10 +34,10 @@ build/final.xclbin: build/aie.mlir build/reduce_min.cc.o cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts.txt $(<:%=../%) -${targetname}.exe: test.cpp +${targetname}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}.exe $@ diff --git a/programming_examples/basic/vector_reduce_min/run.lit b/programming_examples/basic/vector_reduce_min/run.lit deleted file mode 100644 index 710a9a02cd..0000000000 --- a/programming_examples/basic/vector_reduce_min/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../aie_kernels/aie2/reduce_min.cc -o reduce_min.cc.o -// RUN: %python %S/aie2.py npu 0 | aie-opt -cse -canonicalize -o ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++ %S/test.cpp -o test.exe -std=c++23 -Wall -I%S/../../../runtime_lib/test_lib %S/../../../runtime_lib/test_lib/test_utils.cpp %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/vector_reduce_min/run_makefile.lit b/programming_examples/basic/vector_reduce_min/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_reduce_min/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_scalar_add/Makefile b/programming_examples/basic/vector_scalar_add/Makefile index 463b63532b..1f61250c87 100644 --- a/programming_examples/basic/vector_scalar_add/Makefile +++ b/programming_examples/basic/vector_scalar_add/Makefile @@ -6,13 +6,15 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) + +include ${srcdir}/../../makefile-common all: build/final.xclbin targetname = vectorScalarAdd -build/aie.mlir: aie2.py +build/aie.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< > $@ @@ -21,10 +23,10 @@ build/final.xclbin: build/aie.mlir cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ --xclbin-name=${@F} --npu-insts-name=insts.txt ${ ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! - diff --git a/programming_examples/basic/vector_scalar_add/run_makefile.lit b/programming_examples/basic/vector_scalar_add/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_scalar_add/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_scalar_mul/Makefile b/programming_examples/basic/vector_scalar_mul/Makefile index ee55a2ae2b..f4388f5a6b 100644 --- a/programming_examples/basic/vector_scalar_mul/Makefile +++ b/programming_examples/basic/vector_scalar_mul/Makefile @@ -6,9 +6,11 @@ # ##===----------------------------------------------------------------------===## -include ../../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) -VPATH := ../../../aie_kernels/aie2 +include ${srcdir}/../../makefile-common + +VPATH := ${srcdir}/../../../aie_kernels/aie2 targetname = vectorScalar data_size = 4096 @@ -20,9 +22,9 @@ kristof: build/insts_${data_size}.txt build/%.o: %.cc mkdir -p ${@D} - cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $(<:%=../%) -o ${@F} + cd ${@D} && xchesscc_wrapper ${CHESSCCWRAP2_FLAGS} -c $< -o ${@F} -build/aie_${data_size}.mlir: aie2.py +build/aie_${data_size}.mlir: ${srcdir}/aie2.py mkdir -p ${@D} python3 $< ${data_size} 0 > $@ @@ -41,10 +43,10 @@ build/final_trace_${data_size}.xclbin: build/aie_trace_${data_size}.mlir build/s cd ${@D} && aiecc.py --aie-generate-cdo --no-compile-host --xclbin-name=${@F} \ --aie-generate-npu --npu-insts-name=insts_${data_size}.txt $(<:%=../%) -${targetname}_${data_size}.exe: test.cpp +${targetname}_${data_size}.exe: ${srcdir}/test.cpp rm -rf _build mkdir -p _build - cd _build && ${powershell} cmake .. -DTARGET_NAME=${targetname}_${data_size} -DVECTORSCALARMUL_SIZE=${data_size} + cd _build && ${powershell} cmake ${srcdir} -DTARGET_NAME=${targetname}_${data_size} -DVECTORSCALARMUL_SIZE=${data_size} cd _build && ${powershell} cmake --build . --config Release ifeq "${powershell}" "powershell.exe" cp _build/${targetname}_${data_size}.exe $@ @@ -56,14 +58,14 @@ run: ${targetname}_${data_size}.exe build/final_${data_size}.xclbin build/insts_ ${powershell} ./$< -x build/final_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE run_py: build/final_${data_size}.xclbin build/insts_${data_size}.txt - ${powershell} python3 test.py -x build/final_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -s ${data_size} + ${powershell} python3 ${srcdir}/test.py -x build/final_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -s ${data_size} trace: ${targetname}_${data_size}.exe build/final_trace_${data_size}.xclbin build/insts_${data_size}.txt ${powershell} ./$< -x build/final_trace_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -t ${trace_size} ../../utils/parse_trace.py --filename trace.txt --mlir build/aie_trace_${data_size}.mlir --colshift 1 > trace_vs.json trace_py: build/final_trace_${data_size}.xclbin build/insts_${data_size}.txt - ${powershell} python3 test.py -x build/final_trace_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -t ${trace_size} -s ${data_size} + ${powershell} python3 ${srcdir}/test.py -x build/final_trace_${data_size}.xclbin -i build/insts_${data_size}.txt -k MLIR_AIE -t ${trace_size} -s ${data_size} ../../utils/parse_trace.py --filename trace.txt --mlir build/aie_trace_${data_size}.mlir --colshift 1 > trace_vs.json diff --git a/programming_examples/basic/vector_scalar_mul/run.lit b/programming_examples/basic/vector_scalar_mul/run.lit deleted file mode 100644 index fd55fdb97d..0000000000 --- a/programming_examples/basic/vector_scalar_mul/run.lit +++ /dev/null @@ -1,12 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai, chess -// -// RUN: xchesscc_wrapper aie2 -I %aietools/include -c %S/../../../aie_kernels/aie2/scale.cc -o ./scale.o -// RUN: %python %S/aie2.py 4096 0 > ./aie.mlir -// RUN: %python aiecc.py --xbridge --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: g++ %S/test.cpp -o test.exe -std=c++23 -Wall -DVECTORSCALARMUL_SIZE=4096 -I%S/../../../runtime_lib/test_lib %S/../../../runtime_lib/test_lib/test_utils.cpp %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// RUN: %run_on_npu %python %S/test.py -x aie.xclbin -i insts.txt -k MLIR_AIE -s 4096 | FileCheck %s -// CHECK: PASS! diff --git a/programming_examples/basic/vector_scalar_mul/run_makefile.lit b/programming_examples/basic/vector_scalar_mul/run_makefile.lit new file mode 100644 index 0000000000..eb23a0be03 --- /dev/null +++ b/programming_examples/basic/vector_scalar_mul/run_makefile.lit @@ -0,0 +1,10 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// RUN: %run_on_npu make -f %S/Makefile run_py | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_sum/Makefile b/programming_examples/basic/vector_sum/Makefile deleted file mode 100755 index e9c2016543..0000000000 --- a/programming_examples/basic/vector_sum/Makefile +++ /dev/null @@ -1,63 +0,0 @@ -##===- Makefile -----------------------------------------------------------===## -# -# This file 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 -# -##===----------------------------------------------------------------------===## - -include ../makefile-common - -ACDC_AIE = $(dir $(shell which aie-opt))/.. - -SHELL := /bin/bash - -targetname = vectorSum -devicename = npu -col = 0 - -all: build/final.xclbin - -build/final.xclbin: build/aie.mlir - mkdir -p ${@D} - cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ - --xclbin-name=${@F} --npu-insts-name=insts.txt ${ $@ - -# Changing variables when we target VCK5000 -vck5000: devicename=xcvc1902 -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 - -run: ${targetname}.exe build/final.xclbin build/insts.txt - ${powershell} ./$< -x build/final.xclbin -i build/insts.txt -k MLIR_AIE - -clean: - rm -rf build _build inst ${targetname}.exe aie.mlir.prj core_* test.elf diff --git a/programming_examples/basic/vector_sum/README.md b/programming_examples/basic/vector_sum/README.md deleted file mode 100644 index 60d9715528..0000000000 --- a/programming_examples/basic/vector_sum/README.md +++ /dev/null @@ -1,26 +0,0 @@ - - -# Vector sum - -Single tile traverses through a vector in memory and returns the sum of each value in the vector. The tile that performs the operation is tile (`col`, 2) and the data is read from and written to external memory through Shim tile (`col`, 0). A buffer in tile (`col`, 2) is used to store the temporary maximum value during processing, which is then pushed through an object FIFO to the Shim tile when processing is complete. This reference design can be run on either a RyzenAI NPU or a VCK5000. The value of `col` is dependent on whether the application is targetting NPU or VCK5000. - -To compile and run the design for NPU: -``` -make -make run -``` - -To compile and run the design for VCK5000: -``` -make vck5000 -./test.elf -``` - diff --git a/programming_examples/basic/vector_sum/aie2.py b/programming_examples/basic/vector_sum/aie2.py deleted file mode 100755 index 8073833962..0000000000 --- a/programming_examples/basic/vector_sum/aie2.py +++ /dev/null @@ -1,87 +0,0 @@ -# -# 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 -# -# (c) Copyright 2023 AMD Inc. - -import sys - -from aie.dialects.aie import * -from aie.dialects.aiex import * -from aie.dialects.scf import * -from aie.extras.context import mlir_mod_ctx -from aie.extras.dialects.ext import memref, arith - -import sys - - -def my_vector_sum(): - N = 64 - - buffer_depth = 2 - - with mlir_mod_ctx() as ctx: - - if len(sys.argv) != 3: - raise ValueError("[ERROR] Need 2 command line arguments (Device name, Col)") - - if sys.argv[1] == "npu": - dev = AIEDevice.npu - elif sys.argv[1] == "xcvc1902": - dev = AIEDevice.xcvc1902 - else: - raise ValueError("[ERROR] Device name {} is unknown".format(sys.argv[1])) - - @device(dev) - def device_body(): - memRef_ty = T.memref(N, T.i32()) - - # AIE Core Function declarations - - # Tile declarations - ShimTile = tile(int(sys.argv[2]), 0) - ComputeTile2 = tile(int(sys.argv[2]), 2) - - # AIE-array data movement with object fifos - of_in = object_fifo("in", ShimTile, ComputeTile2, buffer_depth, memRef_ty) - of_out = object_fifo("out", ComputeTile2, ShimTile, buffer_depth, memRef_ty) - - # Set up compute tiles - - # Compute tile 2 - @core(ComputeTile2) - def core_body(): - sum_val = memref.alloc(1, T.i32()) - memref.store(arith.constant(0, T.i32()), sum_val, [0]) - # Effective while(1) - for _ in for_(sys.maxsize): - # Number of sub-vector "tile" iterations - elem_in = of_in.acquire(ObjectFifoPort.Consume, 1) - elem_out = of_out.acquire(ObjectFifoPort.Produce, 1) - for i in for_(N): - v0 = memref.load(elem_in, [i]) - v1 = memref.load(sum_val, [0]) - v2 = arith.addi(v1, v0) - memref.store(v2, sum_val, [0]) - yield_([]) - - v3 = memref.load(sum_val, [0]) - memref.store(v3, elem_out, [0]) - of_in.release(ObjectFifoPort.Consume, 1) - of_out.release(ObjectFifoPort.Produce, 1) - yield_([]) - - # To/from AIE-array data movement - tensor_ty = T.memref(N, T.i32()) - - @FuncOp.from_py_func(tensor_ty, tensor_ty, tensor_ty) - def sequence(A, B, C): - npu_dma_memcpy_nd(metadata="out", bd_id=0, mem=C, sizes=[1, 1, 1, 1]) - npu_dma_memcpy_nd(metadata="in", bd_id=1, mem=A, sizes=[1, 1, 1, N]) - npu_sync(column=0, row=0, direction=0, channel=0) - - print(ctx.module) - - -my_vector_sum() diff --git a/programming_examples/basic/vector_sum/run.lit b/programming_examples/basic/vector_sum/run.lit deleted file mode 100644 index 1922c01828..0000000000 --- a/programming_examples/basic/vector_sum/run.lit +++ /dev/null @@ -1,11 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: ryzen_ai -// -// RUN: %python %S/aie2.py npu 0 > ./aie.mlir -// RUN: %python aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host --xclbin-name=aie.xclbin --npu-insts-name=insts.txt ./aie.mlir -// RUN: clang %S/test.cpp -o test.exe -std=c++11 -Wall %xrt_flags -lrt -lstdc++ -lboost_program_options -lboost_filesystem -// RUN: %run_on_npu ./test.exe -x aie.xclbin -k MLIR_AIE -i insts.txt | FileCheck %s -// CHECK: PASS! - diff --git a/programming_examples/basic/vector_sum/run_vck5000.lit b/programming_examples/basic/vector_sum/run_vck5000.lit deleted file mode 100644 index 83b350c50d..0000000000 --- a/programming_examples/basic/vector_sum/run_vck5000.lit +++ /dev/null @@ -1,9 +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 - diff --git a/programming_examples/basic/vector_sum/test.cpp b/programming_examples/basic/vector_sum/test.cpp deleted file mode 100644 index b87b46b888..0000000000 --- a/programming_examples/basic/vector_sum/test.cpp +++ /dev/null @@ -1,192 +0,0 @@ -//===- test.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) 2023, Advanced Micro Devices, Inc. -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include -#include -#include -#include - -#include "xrt/xrt_bo.h" -#include "xrt/xrt_device.h" -#include "xrt/xrt_kernel.h" - -constexpr int IN_SIZE = 64; -constexpr int OUT_SIZE = 1; - -namespace po = boost::program_options; - -void check_arg_file_exists(po::variables_map &vm_in, std::string name) { - if (!vm_in.count(name)) { - throw std::runtime_error("Error: no " + name + " file was provided\n"); - } else { - std::ifstream test(vm_in[name].as()); - if (!test) { - throw std::runtime_error("The " + name + " file " + - vm_in[name].as() + - " does not exist.\n"); - } - } -} - -std::vector load_instr_sequence(std::string instr_path) { - std::ifstream instr_file(instr_path); - std::string line; - std::vector instr_v; - while (std::getline(instr_file, line)) { - std::istringstream iss(line); - uint32_t a; - if (!(iss >> std::hex >> a)) { - throw std::runtime_error("Unable to parse instruction file\n"); - } - instr_v.push_back(a); - } - return instr_v; -} - -int main(int argc, const char *argv[]) { - - // Program arguments parsing - po::options_description desc("Allowed options"); - desc.add_options()("help,h", "produce help message")( - "xclbin,x", po::value()->required(), - "the input xclbin path")( - "kernel,k", po::value()->required(), - "the kernel name in the XCLBIN (for instance PP_PRE_FD)")( - "verbosity,v", po::value()->default_value(0), - "the verbosity of the output")( - "instr,i", po::value()->required(), - "path of file containing userspace instructions to be sent to the LX6"); - po::variables_map vm; - - try { - po::store(po::parse_command_line(argc, argv, desc), vm); - po::notify(vm); - - if (vm.count("help")) { - std::cout << desc << "\n"; - return 1; - } - } catch (const std::exception &ex) { - std::cerr << ex.what() << "\n\n"; - std::cerr << "Usage:\n" << desc << "\n"; - return 1; - } - - check_arg_file_exists(vm, "xclbin"); - check_arg_file_exists(vm, "instr"); - - std::vector instr_v = - load_instr_sequence(vm["instr"].as()); - - int verbosity = vm["verbosity"].as(); - if (verbosity >= 1) - std::cout << "Sequence instr count: " << instr_v.size() << "\n"; - - // Start the XRT test code - // Get a device handle - unsigned int device_index = 0; - auto device = xrt::device(device_index); - - // Load the xclbin - if (verbosity >= 1) - std::cout << "Loading xclbin: " << vm["xclbin"].as() << "\n"; - auto xclbin = xrt::xclbin(vm["xclbin"].as()); - - if (verbosity >= 1) - std::cout << "Kernel opcode: " << vm["kernel"].as() << "\n"; - std::string Node = vm["kernel"].as(); - - // Get the kernel from the xclbin - auto xkernels = xclbin.get_kernels(); - auto xkernel = *std::find_if(xkernels.begin(), xkernels.end(), - [Node](xrt::xclbin::kernel &k) { - auto name = k.get_name(); - std::cout << "Name: " << name << std::endl; - return name.rfind(Node, 0) == 0; - }); - auto kernelName = xkernel.get_name(); - - if (verbosity >= 1) - std::cout << "Registering xclbin: " << vm["xclbin"].as() - << "\n"; - - device.register_xclbin(xclbin); - - // get a hardware context - if (verbosity >= 1) - std::cout << "Getting hardware context.\n"; - xrt::hw_context context(device, xclbin.get_uuid()); - - // get a kernel handle - if (verbosity >= 1) - std::cout << "Getting handle to kernel:" << kernelName << "\n"; - auto kernel = xrt::kernel(context, kernelName); - - auto bo_instr = xrt::bo(device, instr_v.size() * sizeof(int), - XCL_BO_FLAGS_CACHEABLE, kernel.group_id(0)); - auto bo_inA = xrt::bo(device, IN_SIZE * sizeof(int32_t), - XRT_BO_FLAGS_HOST_ONLY, kernel.group_id(2)); - auto bo_inB = xrt::bo(device, IN_SIZE * sizeof(int32_t), - XRT_BO_FLAGS_HOST_ONLY, kernel.group_id(3)); - auto bo_out = xrt::bo(device, OUT_SIZE * sizeof(int32_t), - XRT_BO_FLAGS_HOST_ONLY, kernel.group_id(4)); - - if (verbosity >= 1) - std::cout << "Writing data into buffer objects.\n"; - - int32_t *bufInA = bo_inA.map(); - std::vector srcVecA; - for (int i = 0; i < IN_SIZE; i++) - srcVecA.push_back(i + 1); - memcpy(bufInA, srcVecA.data(), (srcVecA.size() * sizeof(uint32_t))); - - bufInA[IN_SIZE / 2] = 654321; - bufInA[IN_SIZE - 1] = 100; - - void *bufInstr = bo_instr.map(); - memcpy(bufInstr, instr_v.data(), instr_v.size() * sizeof(int)); - - bo_instr.sync(XCL_BO_SYNC_BO_TO_DEVICE); - bo_inA.sync(XCL_BO_SYNC_BO_TO_DEVICE); - - if (verbosity >= 1) - std::cout << "Running Kernel.\n"; - auto run = kernel(bo_instr, instr_v.size(), bo_inA, bo_inB, bo_out); - run.wait(); - - bo_out.sync(XCL_BO_SYNC_BO_FROM_DEVICE); - - uint32_t *bufOut = bo_out.map(); - - int errors = 0; - - uint32_t sum_val = 0; - for (uint32_t i = 0; i < IN_SIZE; i++) { - sum_val += *(bufInA + i); - } - - if (*bufOut != sum_val) { - std::cout << "[ERROR] Maximum value is " << sum_val - << " but kernel returned " << *bufOut << "\n"; - errors++; - } - - if (!errors) { - std::cout << "\nPASS!\n\n"; - return 0; - } else { - std::cout << "\nfailed.\n\n"; - return 1; - } -} diff --git a/programming_examples/basic/vector_sum/test_vck5000.cpp b/programming_examples/basic/vector_sum/test_vck5000.cpp deleted file mode 100644 index aad99cf82d..0000000000 --- a/programming_examples/basic/vector_sum/test_vck5000.cpp +++ /dev/null @@ -1,138 +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); - -#define DMA_COUNT 64 - - // 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; - } - - // 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 sum_val = 0; - for (uint32_t i = 0; i < DMA_COUNT; i++) { - sum_val += in_a[i]; - } - - if (*out != sum_val) { - errors++; - printf("[ERROR] Maximum value is %d but kernel returned %d\n", sum_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_sum/CMakeLists.txt b/programming_examples/basic/vector_vector_add/CMakeLists.txt similarity index 95% rename from programming_examples/basic/vector_sum/CMakeLists.txt rename to programming_examples/basic/vector_vector_add/CMakeLists.txt index 5e637b4d7d..20f5d8a4a3 100644 --- a/programming_examples/basic/vector_sum/CMakeLists.txt +++ b/programming_examples/basic/vector_vector_add/CMakeLists.txt @@ -13,7 +13,7 @@ # cmake needs this line cmake_minimum_required(VERSION 3.1) -set(CMAKE_CXX_STANDARD 23) +set(CMAKE_CXX_STANDARD 23) set(CMAKE_CXX_STANDARD_REQUIRED YES) find_program(WSL NAMES powershell.exe) @@ -35,7 +35,6 @@ set(TARGET_NAME test CACHE STRING "Target to be built") SET (ProjectName ${TARGET_NAME}) SET (currentTarget ${TARGET_NAME}) - if ( WSL ) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY_RELEASE ${CMAKE_BINARY_DIR}) endif () @@ -50,8 +49,6 @@ add_executable(${currentTarget} test.cpp ) -set_property(TARGET ${currentTarget} PROPERTY CXX_STANDARD 23) - target_compile_definitions(${currentTarget} PUBLIC DISABLE_ABI_CHECK=1) target_include_directories (${currentTarget} PUBLIC diff --git a/programming_examples/basic/vector_add/Makefile b/programming_examples/basic/vector_vector_add/Makefile similarity index 88% rename from programming_examples/basic/vector_add/Makefile rename to programming_examples/basic/vector_vector_add/Makefile index 61133a555b..9d732bb0e3 100755 --- a/programming_examples/basic/vector_add/Makefile +++ b/programming_examples/basic/vector_vector_add/Makefile @@ -6,11 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) -ACDC_AIE = $(dir $(shell which aie-opt))/.. - -SHELL := /bin/bash +include ${srcdir}/../../makefile-common targetname = vectorAdd devicename = npu @@ -18,15 +16,19 @@ col = 0 all: build/final.xclbin +build/aie.mlir: ${srcdir}/aie2.py + mkdir -p ${@D} + python3 $< ${devicename} ${col} > $@ + build/final.xclbin: build/aie.mlir mkdir -p ${@D} cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ --xclbin-name=${@F} --npu-insts-name=insts.txt ${ $@ - # Changing variables when we target VCK5000 +ACDC_AIE = $(dir $(shell which aie-opt))/.. + vck5000: devicename=xcvc1902 vck5000: col=6 diff --git a/programming_examples/basic/vector_add/README.md b/programming_examples/basic/vector_vector_add/README.md similarity index 95% rename from programming_examples/basic/vector_add/README.md rename to programming_examples/basic/vector_vector_add/README.md index 65cdafefca..2ed5a82605 100644 --- a/programming_examples/basic/vector_add/README.md +++ b/programming_examples/basic/vector_vector_add/README.md @@ -8,7 +8,7 @@ // //===----------------------------------------------------------------------===//--> -# Vector Add +# Vector Vector Add Single tile performs a very simple `+` operations from two vectors loaded into memory. The tile then stores the sum of those two vectors back to external memory. This reference design can be run on either a RyzenAI NPU or a VCK5000. diff --git a/programming_examples/basic/vector_add/aie2.py b/programming_examples/basic/vector_vector_add/aie2.py similarity index 100% rename from programming_examples/basic/vector_add/aie2.py rename to programming_examples/basic/vector_vector_add/aie2.py diff --git a/programming_examples/basic/vector_vector_add/run_makefile.lit b/programming_examples/basic/vector_vector_add/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_vector_add/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_add/run_vck5000.lit b/programming_examples/basic/vector_vector_add/run_vck5000.lit similarity index 100% rename from programming_examples/basic/vector_add/run_vck5000.lit rename to programming_examples/basic/vector_vector_add/run_vck5000.lit diff --git a/programming_examples/basic/vector_add/test.cpp b/programming_examples/basic/vector_vector_add/test.cpp similarity index 97% rename from programming_examples/basic/vector_add/test.cpp rename to programming_examples/basic/vector_vector_add/test.cpp index c784db343d..550915c698 100644 --- a/programming_examples/basic/vector_add/test.cpp +++ b/programming_examples/basic/vector_vector_add/test.cpp @@ -181,8 +181,9 @@ int main(int argc, const char *argv[]) { << std::endl; errors++; } else { - std::cout << "Correct output " << *(bufOut + i) - << " == " << *(bufInA + i) + *(bufInB + i) << std::endl; + if (verbosity > 1) + std::cout << "Correct output " << *(bufOut + i) + << " == " << *(bufInA + i) + *(bufInB + i) << std::endl; } } diff --git a/programming_examples/basic/vector_add/test_vck5000.cpp b/programming_examples/basic/vector_vector_add/test_vck5000.cpp similarity index 100% rename from programming_examples/basic/vector_add/test_vck5000.cpp rename to programming_examples/basic/vector_vector_add/test_vck5000.cpp diff --git a/programming_examples/basic/vector_vector_mul/CMakeLists.txt b/programming_examples/basic/vector_vector_mul/CMakeLists.txt new file mode 100644 index 0000000000..20f5d8a4a3 --- /dev/null +++ b/programming_examples/basic/vector_vector_mul/CMakeLists.txt @@ -0,0 +1,75 @@ +# 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 +# +# (c) Copyright 2023 Advanced Micro Devices, Inc. + +# parameters +# -DBOOST_ROOT: Path to Boost install +# -DXRT_INC_DIR: Full path to src/runtime_src/core/include in XRT cloned repo +# -DXRT_LIB_DIR: Path to xrt_coreutil.lib +# -DTARGET_NAME: Target name to be built + +# cmake needs this line +cmake_minimum_required(VERSION 3.1) + +set(CMAKE_CXX_STANDARD 23) +set(CMAKE_CXX_STANDARD_REQUIRED YES) + +find_program(WSL NAMES powershell.exe) + +if (NOT WSL) + set(CMAKE_C_COMPILER gcc-13) + set(CMAKE_CXX_COMPILER g++-13) + set(BOOST_ROOT /usr/include/boost CACHE STRING "Path to Boost install") + set(XRT_INC_DIR /opt/xilinx/xrt/include CACHE STRING "Path to XRT cloned repo") + set(XRT_LIB_DIR /opt/xilinx/xrt/lib CACHE STRING "Path to xrt_coreutil.lib") +else() + set(BOOST_ROOT C:/Technical/thirdParty/boost_1_83_0 CACHE STRING "Path to Boost install") + set(XRT_INC_DIR C:/Technical/XRT/src/runtime_src/core/include CACHE STRING "Path to XRT cloned repo") + set(XRT_LIB_DIR C:/Technical/xrtNPUfromDLL CACHE STRING "Path to xrt_coreutil.lib") +endif() + +set(TARGET_NAME test CACHE STRING "Target to be built") + +SET (ProjectName ${TARGET_NAME}) +SET (currentTarget ${TARGET_NAME}) + +if ( WSL ) + set(CMAKE_RUNTIME_OUTPUT_DIRECTORY_RELEASE ${CMAKE_BINARY_DIR}) +endif () + +project(${ProjectName}) + +# Find packages +find_package(Boost REQUIRED) + +add_executable(${currentTarget} + ${CMAKE_CURRENT_SOURCE_DIR}/../../../runtime_lib/test_lib/test_utils.cpp + test.cpp +) + +target_compile_definitions(${currentTarget} PUBLIC DISABLE_ABI_CHECK=1) + +target_include_directories (${currentTarget} PUBLIC + ${XRT_INC_DIR} + ${Boost_INCLUDE_DIRS} + ${CMAKE_CURRENT_SOURCE_DIR}/../../../runtime_lib/test_lib +) + +target_link_directories(${currentTarget} PUBLIC + ${XRT_LIB_DIR} + ${Boost_LIBRARY_DIRS} +) + +if (NOT WSL) + target_link_libraries(${currentTarget} PUBLIC + xrt_coreutil + boost_program_options + boost_filesystem + ) +else() + target_link_libraries(${currentTarget} PUBLIC + xrt_coreutil + ) +endif() diff --git a/programming_examples/basic/vector_mult/Makefile b/programming_examples/basic/vector_vector_mul/Makefile similarity index 87% rename from programming_examples/basic/vector_mult/Makefile rename to programming_examples/basic/vector_vector_mul/Makefile index bc07e3d05b..6697a884ca 100755 --- a/programming_examples/basic/vector_mult/Makefile +++ b/programming_examples/basic/vector_vector_mul/Makefile @@ -6,11 +6,9 @@ # ##===----------------------------------------------------------------------===## -include ../makefile-common +srcdir := $(shell dirname $(realpath $(firstword $(MAKEFILE_LIST)))) -ACDC_AIE = $(dir $(shell which aie-opt))/.. - -SHELL := /bin/bash +include ${srcdir}/../../makefile-common targetname = vectorMult devicename = npu @@ -18,15 +16,19 @@ col = 0 all: build/final.xclbin +build/aie.mlir: ${srcdir}/aie2.py + mkdir -p ${@D} + python3 $< ${devicename} ${col} > $@ + build/final.xclbin: build/aie.mlir mkdir -p ${@D} cd ${@D} && aiecc.py --aie-generate-cdo --aie-generate-npu --no-compile-host \ --xclbin-name=${@F} --npu-insts-name=insts.txt ${ $@ + # Changing variables when we target VCK5000 +ACDC_AIE = $(dir $(shell which aie-opt))/.. + vck5000: devicename=xcvc1902 vck5000: col=6 diff --git a/programming_examples/basic/vector_mult/README.md b/programming_examples/basic/vector_vector_mul/README.md similarity index 94% rename from programming_examples/basic/vector_mult/README.md rename to programming_examples/basic/vector_vector_mul/README.md index 3abe2b9999..54ab0bd4e1 100644 --- a/programming_examples/basic/vector_mult/README.md +++ b/programming_examples/basic/vector_vector_mul/README.md @@ -8,7 +8,7 @@ // //===----------------------------------------------------------------------===//--> -# Vector Multiplication +# Vector Vector Multiplication Single tile performs a very simple `*` operations from two vectors loaded into memory. The tile then stores the element wise multiplication of those two vectors back to external memory. This reference design can be run on either a RyzenAI NPU or a VCK5000. diff --git a/programming_examples/basic/vector_mult/aie2.py b/programming_examples/basic/vector_vector_mul/aie2.py similarity index 100% rename from programming_examples/basic/vector_mult/aie2.py rename to programming_examples/basic/vector_vector_mul/aie2.py diff --git a/programming_examples/basic/vector_vector_mul/run_makefile.lit b/programming_examples/basic/vector_vector_mul/run_makefile.lit new file mode 100644 index 0000000000..6875524001 --- /dev/null +++ b/programming_examples/basic/vector_vector_mul/run_makefile.lit @@ -0,0 +1,9 @@ +// (c) Copyright 2024 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// REQUIRES: ryzen_ai, chess +// +// RUN: make -f %S/Makefile clean +// RUN: make -f %S/Makefile +// RUN: %run_on_npu make -f %S/Makefile run | FileCheck %s +// CHECK: PASS! diff --git a/programming_examples/basic/vector_mult/run_vck5000.lit b/programming_examples/basic/vector_vector_mul/run_vck5000.lit similarity index 100% rename from programming_examples/basic/vector_mult/run_vck5000.lit rename to programming_examples/basic/vector_vector_mul/run_vck5000.lit diff --git a/programming_examples/basic/vector_mult/test.cpp b/programming_examples/basic/vector_vector_mul/test.cpp similarity index 97% rename from programming_examples/basic/vector_mult/test.cpp rename to programming_examples/basic/vector_vector_mul/test.cpp index 3e9990f0ee..1ba8c8159e 100644 --- a/programming_examples/basic/vector_mult/test.cpp +++ b/programming_examples/basic/vector_vector_mul/test.cpp @@ -182,8 +182,9 @@ int main(int argc, const char *argv[]) { << std::endl; errors++; } else { - std::cout << "Correct output " << *(bufOut + i) - << " == " << *(bufInA + i) * *(bufInB + i) << std::endl; + if (verbosity >= 1) + std::cout << "Correct output " << *(bufOut + i) + << " == " << *(bufInA + i) * *(bufInB + i) << std::endl; } } diff --git a/programming_examples/basic/vector_mult/test_vck5000.cpp b/programming_examples/basic/vector_vector_mul/test_vck5000.cpp similarity index 100% rename from programming_examples/basic/vector_mult/test_vck5000.cpp rename to programming_examples/basic/vector_vector_mul/test_vck5000.cpp diff --git a/programming_examples/makefile-common b/programming_examples/makefile-common index 42565c463f..a95c675d68 100644 --- a/programming_examples/makefile-common +++ b/programming_examples/makefile-common @@ -1,7 +1,3 @@ -# Contains common definitions used across the Makefiles of npu-xrt tests. -REPO_ROOT ?= $(shell realpath $(dir $(shell which aie-opt))/../../..) -INSTALL_ROOT ?= $(shell realpath $(dir $(shell which aie-opt))/..) - # VITIS related variables VITIS_ROOT ?= $(shell realpath $(dir $(shell which vitis))/../) VITIS_AIETOOLS_DIR ?= ${VITIS_ROOT}/aietools