From 0d8585868d430e8796a57ed9886747f8ea20bf8a Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Wed, 19 Nov 2025 10:56:32 +0800 Subject: [PATCH 1/9] modify samples for cpu/npu unify --- .../MatmulInvocationNeo/CMakeLists.txt | 2 +- .../MatmulInvocationNeo/cmake/cpu_lib.cmake | 5 +++- .../MatmulInvocationNeo/main.cpp | 27 ----------------- .../AddKernelInvocationNeo/CMakeLists.txt | 3 +- .../cmake/cpu_lib.cmake | 4 ++- .../AddKernelInvocationNeo/main.cpp | 27 ++--------------- .../CMakeLists.txt | 3 +- .../cmake/cpu_lib.cmake | 4 ++- .../AddKernelInvocationTilingNeo/main.cpp | 29 +------------------ .../5_addn_kernellaunch/CMakeLists.txt | 2 +- .../5_addn_kernellaunch/cmake/cpu_lib.cmake | 4 ++- .../5_addn_kernellaunch/main.cpp | 26 ----------------- 12 files changed, 22 insertions(+), 114 deletions(-) diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt index c85dad2b3..254015b8e 100644 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt @@ -44,7 +44,7 @@ target_compile_definitions(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> + $:ascendc_acl_stub>> ascendc_kernels_${RUN_MODE} tiling_api register diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake index acb98ec90..7fed07d24 100644 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake +++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake @@ -3,7 +3,10 @@ if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) endif() find_package(tikicpulib REQUIRED) -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) +set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +include(${CPU_CMAKE_PATH}/ascendc.cmake) + +ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) target_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE $<$>:CUSTOM_ASCEND310P> diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp index 1fb5a80d8..44ab81f74 100644 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp +++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp @@ -10,13 +10,8 @@ #include "data_utils.h" #include "kernel_tiling/kernel_tiling.h" #include "tiling/platform/platform_ascendc.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_matmul_custom.h" -#else -#include "tikicpulib.h" -extern "C" void matmul_custom(uint8_t *a, uint8_t *b, uint8_t *c, uint8_t *workspace, uint8_t *tiling); -#endif extern void GenerateTiling(const char *socVersion, uint8_t *tilingBuf); int32_t main(int32_t argc, char *argv[]) @@ -39,27 +34,6 @@ int32_t main(int32_t argc, char *argv[]) uint32_t blockDim = (reinterpret_cast(tilingBuf)->usedCoreNum + 1) / 2; #endif -#ifdef ASCENDC_CPU_DEBUG - uint8_t *a = (uint8_t *)AscendC::GmAlloc(aFileSize); - uint8_t *b = (uint8_t *)AscendC::GmAlloc(bFileSize); - uint8_t *c = (uint8_t *)AscendC::GmAlloc(cFileSize); - uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(workspaceSize); - uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tilingFileSize); - - ReadFile("./input/x1_gm.bin", aFileSize, a, aFileSize); - ReadFile("./input/x2_gm.bin", bFileSize, b, bFileSize); - memcpy_s(tiling, tilingFileSize, tilingBuf, tilingFileSize); - - ICPU_RUN_KF(matmul_custom, blockDim, a, b, c, workspace, tiling); - - WriteFile("./output/output.bin", c, cFileSize); - - AscendC::GmFree((void *)a); - AscendC::GmFree((void *)b); - AscendC::GmFree((void *)c); - AscendC::GmFree((void *)workspace); - AscendC::GmFree((void *)tiling); -#else CHECK_ACL(aclInit(nullptr)); int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); @@ -115,7 +89,6 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif free(tilingBuf); return 0; } \ No newline at end of file diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt index 1e4d6de99..421f8c91d 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt @@ -33,8 +33,9 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> ascendc_kernels_${RUN_MODE} + $:ascendc_acl_stub>> + # ascendcl ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake index 5362c8b5a..3c875de99 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake @@ -2,8 +2,10 @@ if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) endif() find_package(tikicpulib REQUIRED) +set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +include(${CPU_CMAKE_PATH}/ascendc.cmake) -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) +ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/main.cpp b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/main.cpp index d3d8fea33..bb62adc8d 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/main.cpp +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/main.cpp @@ -8,13 +8,8 @@ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. */ #include "data_utils.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" -extern void add_custom_do(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, uint8_t *z); -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z); -#endif +#include "aclrtlaunch_add_custom.h" int32_t main(int32_t argc, char *argv[]) { @@ -22,23 +17,6 @@ int32_t main(int32_t argc, char *argv[]) size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); -#ifdef ASCENDC_CPU_DEBUG - uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *y = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *z = (uint8_t *)AscendC::GmAlloc(outputByteSize); - - ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize); - ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize); - - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug - - WriteFile("./output/output_z.bin", z, outputByteSize); - - AscendC::GmFree((void *)x); - AscendC::GmFree((void *)y); - AscendC::GmFree((void *)z); -#else CHECK_ACL(aclInit(nullptr)); int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); @@ -61,7 +39,7 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); - add_custom_do(blockDim, stream, xDevice, yDevice, zDevice); + ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, xDevice, yDevice, zDevice); CHECK_ACL(aclrtSynchronizeStream(stream)); CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST)); @@ -77,6 +55,5 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif return 0; } diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt index 1e4d6de99..421f8c91d 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt @@ -33,8 +33,9 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> ascendc_kernels_${RUN_MODE} + $:ascendc_acl_stub>> + # ascendcl ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake index 5362c8b5a..3c875de99 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake @@ -2,8 +2,10 @@ if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) endif() find_package(tikicpulib REQUIRED) +set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +include(${CPU_CMAKE_PATH}/ascendc.cmake) -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) +ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/main.cpp b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/main.cpp index 527b07f74..8ef9a00f3 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/main.cpp +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/main.cpp @@ -9,13 +9,9 @@ */ #include "add_custom_tiling.h" #include "data_utils.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_add_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling); -#endif + int32_t main(int32_t argc, char *argv[]) { @@ -24,28 +20,6 @@ int32_t main(int32_t argc, char *argv[]) size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); -#ifdef ASCENDC_CPU_DEBUG - uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tilingSize); - ReadFile("./input/input_tiling.bin", tilingSize, tiling, tilingSize); - uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *y = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *z = (uint8_t *)AscendC::GmAlloc(outputByteSize); - - ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize); - ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize); - - AscendC::SetKernelMode(KernelMode::AIV_MODE); - - ICPU_RUN_KF(add_custom, blockDim, x, y, z, - *reinterpret_cast(tiling)); // use this macro for cpu debug - - WriteFile("./output/output_z.bin", z, outputByteSize); - - AscendC::GmFree((void *)x); - AscendC::GmFree((void *)y); - AscendC::GmFree((void *)z); - AscendC::GmFree((void *)tiling); -#else CHECK_ACL(aclInit(nullptr)); int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); @@ -89,6 +63,5 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif return 0; } diff --git a/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt b/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt index 586f565ec..6c9513332 100644 --- a/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt +++ b/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt @@ -31,7 +31,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> + $:ascendc_acl_stub>> ascendc_kernels_${RUN_MODE} ) diff --git a/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake index 5362c8b5a..3c875de99 100644 --- a/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake +++ b/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake @@ -2,8 +2,10 @@ if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) endif() find_package(tikicpulib REQUIRED) +set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +include(${CPU_CMAKE_PATH}/ascendc.cmake) -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) +ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/5_addn_kernellaunch/main.cpp b/operator/ascendc/0_introduction/5_addn_kernellaunch/main.cpp index 542739954..45190e13b 100644 --- a/operator/ascendc/0_introduction/5_addn_kernellaunch/main.cpp +++ b/operator/ascendc/0_introduction/5_addn_kernellaunch/main.cpp @@ -8,13 +8,8 @@ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. */ #include "data_utils.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_addn_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void addn_custom(GM_ADDR srcList, GM_ADDR z); -#endif int32_t main(int32_t argc, char *argv[]) { @@ -41,26 +36,6 @@ int32_t main(int32_t argc, char *argv[]) uintptr_t dataPtr[TENSOR_DESC_NUM]; } inputDesc; -#ifdef ASCENDC_CPU_DEBUG - uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *y = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *z = (uint8_t *)AscendC::GmAlloc(outputByteSize); - - ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize); - ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize); - - inputDesc = { - (1 + (1 + SHAPE_DIM) * TENSOR_DESC_NUM) * sizeof(uint64_t), {xDesc, yDesc}, {(uintptr_t)x, (uintptr_t)y}}; - - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(addn_custom, blockDim, (uint8_t *)&inputDesc, z); // use this macro for cpu debug - - WriteFile("./output/output_z.bin", z, outputByteSize); - - AscendC::GmFree((void *)x); - AscendC::GmFree((void *)y); - AscendC::GmFree((void *)z); -#else CHECK_ACL(aclInit(nullptr)); int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); @@ -108,6 +83,5 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif return 0; } -- Gitee From 3ba9c57a430a86d79f61aa6c6ffc12c6295eaab1 Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Thu, 20 Nov 2025 01:24:34 +0800 Subject: [PATCH 2/9] update Matmul Neo case --- .../MatmulInvocationNeo/CMakeLists.txt | 30 ++++++++++++++----- .../MatmulInvocationNeo/cmake/cpu_lib.cmake | 15 ---------- .../MatmulInvocationNeo/cmake/npu_lib.cmake | 17 ----------- 3 files changed, 22 insertions(+), 40 deletions(-) delete mode 100644 operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt index 254015b8e..9bcdfb893 100644 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt @@ -20,13 +20,33 @@ file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/matmul_custom.cpp) set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) + +ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE + $<$>:CUSTOM_ASCEND310P> + HAVE_WORKSPACE + HAVE_TILING +) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp ${CMAKE_CURRENT_SOURCE_DIR}/matmul_custom_tiling.cpp @@ -44,13 +64,7 @@ target_compile_definitions(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendc_acl_stub>> ascendc_kernels_${RUN_MODE} - tiling_api - register - platform - ascendalog - dl ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake deleted file mode 100644 index 7fed07d24..000000000 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/cpu_lib.cmake +++ /dev/null @@ -1,15 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -include(${CPU_CMAKE_PATH}/ascendc.cmake) - -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> -) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/npu_lib.cmake deleted file mode 100644 index d9a9fb2ea..000000000 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/cmake/npu_lib.cmake +++ /dev/null @@ -1,17 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) - -ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> - HAVE_WORKSPACE - HAVE_TILING -) -- Gitee From 2b5f5e9cffe03c83fbd6f22d35615983cb9aae0b Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Thu, 20 Nov 2025 01:37:34 +0800 Subject: [PATCH 3/9] update Add Neo case --- .../MatmulInvocationNeo/CMakeLists.txt | 6 ++--- .../AddKernelInvocationNeo/CMakeLists.txt | 23 +++++++++++++++---- .../cmake/cpu_lib.cmake | 11 --------- .../cmake/npu_lib.cmake | 11 --------- 4 files changed, 21 insertions(+), 30 deletions(-) delete mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt index 9bcdfb893..605ba4927 100644 --- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/CMakeLists.txt @@ -39,9 +39,9 @@ endif() include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) # ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) -ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE +ascendc_compile_definitions(ascendc_kernels PRIVATE $<$>:CUSTOM_ASCEND310P> HAVE_WORKSPACE HAVE_TILING @@ -64,7 +64,7 @@ target_compile_definitions(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - ascendc_kernels_${RUN_MODE} + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt index 421f8c91d..0e0942e59 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/add_custom.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_compile_options(ascendc_kernels_bbit PRIVATE @@ -33,9 +48,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - ascendc_kernels_${RUN_MODE} - $:ascendc_acl_stub>> - # ascendcl + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake deleted file mode 100644 index 3c875de99..000000000 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/cpu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) -set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -include(${CPU_CMAKE_PATH}/ascendc.cmake) - -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -- Gitee From c11bcf622e000447520b6407a11b8fd00c11c417 Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Thu, 20 Nov 2025 20:20:57 +0800 Subject: [PATCH 4/9] update Add Tiling Neo case --- .../CMakeLists.txt | 23 +++++++++++++++---- .../cmake/cpu_lib.cmake | 11 --------- .../cmake/npu_lib.cmake | 11 --------- 3 files changed, 18 insertions(+), 27 deletions(-) delete mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt index 421f8c91d..0e0942e59 100644 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/add_custom.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_compile_options(ascendc_kernels_bbit PRIVATE @@ -33,9 +48,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - ascendc_kernels_${RUN_MODE} - $:ascendc_acl_stub>> - # ascendcl + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake deleted file mode 100644 index 3c875de99..000000000 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/cpu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) -set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -include(${CPU_CMAKE_PATH}/ascendc.cmake) - -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTilingNeo/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -- Gitee From 8502828c2f975cbeb53609416bd217fe09655756 Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Fri, 21 Nov 2025 01:20:40 +0800 Subject: [PATCH 5/9] update MatmulLeakyrelu Neo case --- .../MatmulLeakyReluInvocation/CMakeLists.txt | 32 +++++++++++++------ .../cmake/cpu_lib.cmake | 12 ------- .../cmake/npu_lib.cmake | 17 ---------- .../MatmulLeakyReluInvocation/main.cpp | 28 ---------------- .../CMakeLists.txt | 32 +++++++++++++------ .../cmake/cpu_lib.cmake | 12 ------- .../cmake/npu_lib.cmake | 17 ---------- .../MatmulLeakyReluInvocationAsync/main.cpp | 28 ---------------- 8 files changed, 46 insertions(+), 132 deletions(-) delete mode 100644 operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/npu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/CMakeLists.txt b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/CMakeLists.txt index f2915dca3..76f193a73 100644 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/CMakeLists.txt +++ b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/CMakeLists.txt @@ -20,13 +20,33 @@ file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/matmul_leakyrelu_custom.cpp) set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + +ascendc_compile_definitions(ascendc_kernels PRIVATE + $<$>:CUSTOM_ASCEND310P> + HAVE_WORKSPACE + HAVE_TILING +) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp ${CMAKE_CURRENT_SOURCE_DIR}/matmul_leakyrelu_custom_tiling.cpp @@ -44,13 +64,7 @@ target_compile_definitions(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} - tiling_api - register - platform - ascendalog - dl + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/cpu_lib.cmake deleted file mode 100644 index acb98ec90..000000000 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/cpu_lib.cmake +++ /dev/null @@ -1,12 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> -) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/npu_lib.cmake deleted file mode 100644 index d9a9fb2ea..000000000 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/cmake/npu_lib.cmake +++ /dev/null @@ -1,17 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) - -ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> - HAVE_WORKSPACE - HAVE_TILING -) diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/main.cpp b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/main.cpp index f4fe78026..60b97a371 100644 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/main.cpp +++ b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocation/main.cpp @@ -10,13 +10,8 @@ #include "data_utils.h" #include "kernel_tiling/kernel_tiling.h" #include "tiling/platform/platform_ascendc.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_matmul_leakyrelu_custom.h" -#else -#include "tikicpulib.h" -extern "C" void matmul_leakyrelu_custom(uint8_t *, uint8_t *, uint8_t *, uint8_t *, uint8_t *, uint8_t *); -#endif extern void GenerateTiling(const char *socVersion, uint8_t *tilingBuf); int32_t main(int32_t argc, char *argv[]) @@ -39,28 +34,6 @@ int32_t main(int32_t argc, char *argv[]) uint32_t blockDim = 1; #endif -#ifdef ASCENDC_CPU_DEBUG - uint8_t *a = (uint8_t *)AscendC::GmAlloc(aFileSize); - uint8_t *b = (uint8_t *)AscendC::GmAlloc(bFileSize); - uint8_t *bias = (uint8_t *)AscendC::GmAlloc(biasFileSize); - uint8_t *c = (uint8_t *)AscendC::GmAlloc(cFileSize); - uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tilingFileSize); - uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(workspaceSize); - - ReadFile("./input/x1_gm.bin", aFileSize, a, aFileSize); - ReadFile("./input/x2_gm.bin", bFileSize, b, bFileSize); - ReadFile("./input/bias.bin", biasFileSize, bias, biasFileSize); - memcpy_s(tiling, tilingFileSize, tilingBuf, tilingFileSize); - ICPU_RUN_KF(matmul_leakyrelu_custom, blockDim, a, b, bias, c, workspace, tiling); - - WriteFile("./output/output.bin", c, cFileSize); - AscendC::GmFree((void *)a); - AscendC::GmFree((void *)b); - AscendC::GmFree((void *)bias); - AscendC::GmFree((void *)c); - AscendC::GmFree((void *)tiling); - AscendC::GmFree((void *)workspace); -#else CHECK_ACL(aclInit(nullptr)); int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); @@ -125,7 +98,6 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif free(tilingBuf); return 0; } \ No newline at end of file diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/CMakeLists.txt b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/CMakeLists.txt index f2915dca3..76f193a73 100644 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/CMakeLists.txt +++ b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/CMakeLists.txt @@ -20,13 +20,33 @@ file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/matmul_leakyrelu_custom.cpp) set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + +ascendc_compile_definitions(ascendc_kernels PRIVATE + $<$>:CUSTOM_ASCEND310P> + HAVE_WORKSPACE + HAVE_TILING +) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp ${CMAKE_CURRENT_SOURCE_DIR}/matmul_leakyrelu_custom_tiling.cpp @@ -44,13 +64,7 @@ target_compile_definitions(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} - tiling_api - register - platform - ascendalog - dl + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/cpu_lib.cmake deleted file mode 100644 index acb98ec90..000000000 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/cpu_lib.cmake +++ /dev/null @@ -1,12 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> -) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/npu_lib.cmake deleted file mode 100644 index d9a9fb2ea..000000000 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/cmake/npu_lib.cmake +++ /dev/null @@ -1,17 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) - -ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> - HAVE_WORKSPACE - HAVE_TILING -) diff --git a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/main.cpp b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/main.cpp index 6c3fa165b..aeee97b90 100644 --- a/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/main.cpp +++ b/operator/ascendc/0_introduction/13_matmulleakyrelu_kernellaunch/MatmulLeakyReluInvocationAsync/main.cpp @@ -10,13 +10,8 @@ #include "data_utils.h" #include "kernel_tiling/kernel_tiling.h" #include "tiling/platform/platform_ascendc.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_matmul_leakyrelu_custom.h" -#else -#include "tikicpulib.h" -extern "C" void matmul_leakyrelu_custom(uint8_t *, uint8_t *, uint8_t *, uint8_t *, uint8_t *, uint8_t *); -#endif extern void GenerateTiling(const char *socVersion, uint8_t *tilingBuf); int32_t main(int32_t argc, char *argv[]) @@ -38,28 +33,6 @@ int32_t main(int32_t argc, char *argv[]) uint8_t *tilingBuf = (uint8_t *)malloc(tilingFileSize); GenerateTiling(socVersion, tilingBuf); -#ifdef ASCENDC_CPU_DEBUG - uint8_t *a = (uint8_t *)AscendC::GmAlloc(aFileSize); - uint8_t *b = (uint8_t *)AscendC::GmAlloc(bFileSize); - uint8_t *bias = (uint8_t *)AscendC::GmAlloc(biasFileSize); - uint8_t *c = (uint8_t *)AscendC::GmAlloc(cFileSize); - uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tilingFileSize); - uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(workspaceSize); - - ReadFile("./input/x1_gm.bin", aFileSize, a, aFileSize); - ReadFile("./input/x2_gm.bin", bFileSize, b, bFileSize); - ReadFile("./input/bias.bin", biasFileSize, bias, biasFileSize); - memcpy_s(tiling, tilingFileSize, tilingBuf, tilingFileSize); - ICPU_RUN_KF(matmul_leakyrelu_custom, blockDim, a, b, bias, c, workspace, tiling); - - WriteFile("./output/output.bin", c, cFileSize); - AscendC::GmFree((void *)a); - AscendC::GmFree((void *)b); - AscendC::GmFree((void *)bias); - AscendC::GmFree((void *)c); - AscendC::GmFree((void *)tiling); - AscendC::GmFree((void *)workspace); -#else CHECK_ACL(aclInit(nullptr)); int32_t deviceId = 0; CHECK_ACL(aclrtSetDevice(deviceId)); @@ -124,7 +97,6 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif free(tilingBuf); return 0; } \ No newline at end of file -- Gitee From 98277fe5bce6ae4e3cb7af0ef63155b02d12772a Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Fri, 21 Nov 2025 10:36:44 +0800 Subject: [PATCH 6/9] update addn case --- .../5_addn_kernellaunch/CMakeLists.txt | 22 +++++++++++++++---- .../5_addn_kernellaunch/cmake/cpu_lib.cmake | 11 ---------- .../5_addn_kernellaunch/cmake/npu_lib.cmake | 10 --------- 3 files changed, 18 insertions(+), 25 deletions(-) delete mode 100644 operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt b/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt index 6c9513332..ca8e52814 100644 --- a/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt +++ b/operator/ascendc/0_introduction/5_addn_kernellaunch/CMakeLists.txt @@ -16,12 +16,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/addn_custom.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_compile_options(ascendc_kernels_bbit PRIVATE @@ -31,8 +46,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendc_acl_stub>> - ascendc_kernels_${RUN_MODE} + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake deleted file mode 100644 index 3c875de99..000000000 --- a/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/cpu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) -set(CPU_CMAKE_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -include(${CPU_CMAKE_PATH}/ascendc.cmake) - -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/npu_lib.cmake deleted file mode 100644 index 8abf63946..000000000 --- a/operator/ascendc/0_introduction/5_addn_kernellaunch/cmake/npu_lib.cmake +++ /dev/null @@ -1,10 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -- Gitee From 26dfbe4e47d1eb38f032f36506e6da05a49c2835 Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Fri, 21 Nov 2025 11:27:26 +0800 Subject: [PATCH 7/9] add AddTemplate Neo case --- .../CMakeLists.txt | 58 +++++ .../AddKernelInvocationTemplateNeo/README.md | 75 +++++++ .../add_custom.cpp | 91 ++++++++ .../add_custom_struct.h | 22 ++ .../data_utils.h | 203 ++++++++++++++++++ .../AddKernelInvocationTemplateNeo/main.cpp | 73 +++++++ .../AddKernelInvocationTemplateNeo/run.sh | 120 +++++++++++ .../scripts/gen_data.py | 25 +++ .../scripts/verify_result.py | 53 +++++ 9 files changed, 720 insertions(+) create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/CMakeLists.txt create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/README.md create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom.cpp create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom_struct.h create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/data_utils.h create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/main.cpp create mode 100755 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/run.sh create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/gen_data.py create mode 100644 operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/verify_result.py diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/CMakeLists.txt b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/CMakeLists.txt new file mode 100644 index 000000000..0e0942e59 --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/CMakeLists.txt @@ -0,0 +1,58 @@ +cmake_minimum_required(VERSION 3.16) +project(Ascend_c) + +set(RUN_MODE "npu" CACHE STRING "cpu/sim/npu") +set(SOC_VERSION "Ascend310P3" CACHE STRING "system on chip type") +set(ASCEND_CANN_PACKAGE_PATH "/usr/local/Ascend/ascend-toolkit/latest" + CACHE STRING "ASCEND CANN package installation directory" +) +if(NOT CMAKE_BUILD_TYPE) + set(CMAKE_BUILD_TYPE "Debug" CACHE STRING "Build type Release/Debug (default Debug)" FORCE) +endif() +if(CMAKE_INSTALL_PREFIX STREQUAL /usr/local) + set(CMAKE_INSTALL_PREFIX "${CMAKE_CURRENT_LIST_DIR}/out" CACHE STRING "path for install()" FORCE) +endif() + +# ${KERNEL_FILES} are used to compile library, push files written by ascendc in ${KERNEL_FILES}. +# ref to cmake/npu.cmake ascendc_library, cmake/cpu.cmake add_library +file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/add_custom.cpp) + +if("${RUN_MODE}" STREQUAL "cpu") + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) +elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") +else() + message("invalid RUN_MODE: ${RUN_MODE}") +endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + +add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) + +target_compile_options(ascendc_kernels_bbit PRIVATE + $:-g>> + -O2 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0 -Wall -Werror +) + +target_link_libraries(ascendc_kernels_bbit PRIVATE + $,$>:host_intf_pub>> + ascendc_kernels +) + +install(TARGETS ascendc_kernels_bbit + LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} + RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} +) diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/README.md b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/README.md new file mode 100644 index 000000000..4259ec5c7 --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/README.md @@ -0,0 +1,75 @@ +## 目录结构介绍 +``` +├── AddKernelInvocationStructNeo +│ ├── cmake // 编译工程文件 +│ ├── scripts +│ │ ├── gen_data.py // 输入数据和真值数据生成脚本 +│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 +│ ├── add_custom_struct.h // 算子struct实现 +│ ├── add_custom.cpp // 算子kernel实现 +│ ├── CMakeLists.txt // 编译工程文件 +│ ├── data_utils.h // 数据读入写出函数 +│ ├── main.cpp // 主函数，调用算子的应用程序，含CPU域及NPU域调用 +│ └── run.sh // 编译运行算子的脚本 +``` +## 代码实现介绍 +本样例中实现的是固定shape为8*2048的Add算子。 +- kernel实现 + Add算子的数学表达式为： + ``` + z = x + y + ``` + 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用计算接口完成两个输入参数相加，得到最终结果，再搬出到外部存储上。 + + Add算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm和yGm搬运到Local Memory，分别存储在xLocal、yLocal，Compute任务负责对xLocal、yLocal执行加法操作，计算结果存储在zLocal中，CopyOut任务负责将输出数据从zLocal搬运至Global Memory上的输出Tensor zGm中。具体请参考[add_custom.cpp](./add_custom.cpp)。 + +- 调用实现 + 主要通过使用ACLRT_LAUNCH_KERNEL内核调用宏来完成。 + + 应用程序通过ASCENDC_CPU_DEBUG 宏区分代码逻辑运行于CPU侧还是NPU侧。 + +## 运行样例算子 + - 打开样例目录 + 以命令行方式下载样例代码，master分支为例。 + ```bash + cd ${git_clone_path}/samples/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationNeo + ``` + - 配置环境变量 + + 请根据当前环境上CANN开发套件包的[安装方式](https://hiascend.com/document/redirect/CannCommunityInstSoftware)，选择对应配置环境变量的命令。 + - 默认路径，root用户安装CANN软件包 + ```bash + export ASCEND_INSTALL_PATH=/usr/local/Ascend/ascend-toolkit/latest + ``` + - 默认路径，非root用户安装CANN软件包 + ```bash + export ASCEND_INSTALL_PATH=$HOME/Ascend/ascend-toolkit/latest + ``` + - 指定路径install_path，安装CANN软件包 + ```bash + export ASCEND_INSTALL_PATH=${install_path}/ascend-toolkit/latest + ``` + + + + - 样例执行 + + ```bash + bash run.sh -r [RUN_MODE] -v [SOC_VERSION] + ``` + - RUN_MODE：编译方式，可选择CPU调试，NPU仿真，NPU上板。支持参数为[cpu /sim / npu] + - SOC_VERSION：昇腾AI处理器型号，如果无法确定具体的[SOC_VERSION]，则在安装昇腾AI处理器的服务器执行npu-smi info命令进行查询，在查询到的“Name”前增加Ascend信息，例如“Name”对应取值为xxxyy，实际配置的[SOC_VERSION]值为Ascendxxxyy。支持以下产品型号： + - Atlas 训练系列产品 + - Atlas 推理系列产品AI Core + - Atlas A2训练系列产品/Atlas 800I A2推理产品 + - Atlas 200/500 A2推理产品 + + 示例如下，Ascendxxxyy请替换为实际的AI处理器型号。 + ```bash + bash run.sh -r cpu -v Ascendxxxyy + ``` +## 更新说明 +| 时间 | 更新事项 | +| ---------- | ------------ | +| 2024/05/22 | 新增本readme | +| 2024/11/11 | 样例目录调整 | \ No newline at end of file diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom.cpp b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom.cpp new file mode 100644 index 000000000..05043061f --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom.cpp @@ -0,0 +1,91 @@ +/** + * @file add_custom.cpp + * + * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. + */ +#include "kernel_operator.h" +#include "add_custom_struct.h" + +constexpr int32_t TOTAL_LENGTH = 8 * 2048; // total length of data +constexpr int32_t USE_CORE_NUM = 8; // num of core used +constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // length computed of each core +constexpr int32_t TILE_NUM = 8; // split data into 8 tiles for each core +constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue +constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // separate to 2 parts, due to double buffer + +class KernelAdd { +public: + __aicore__ inline KernelAdd() {} + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) + { + xGm.SetGlobalBuffer((__gm__ half *)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); + yGm.SetGlobalBuffer((__gm__ half *)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); + zGm.SetGlobalBuffer((__gm__ half *)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH); + pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half)); + pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half)); + pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half)); + } + __aicore__ inline void Process() + { + int32_t loopCount = TILE_NUM * BUFFER_NUM; + for (int32_t i = 0; i < loopCount; i++) { + CopyIn(i); + Compute(i); + CopyOut(i); + } + } + +private: + __aicore__ inline void CopyIn(int32_t progress) + { + AscendC::LocalTensor xLocal = inQueueX.AllocTensor(); + AscendC::LocalTensor yLocal = inQueueY.AllocTensor(); + AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); + AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); + inQueueX.EnQue(xLocal); + inQueueY.EnQue(yLocal); + } + __aicore__ inline void Compute(int32_t progress) + { + AscendC::LocalTensor xLocal = inQueueX.DeQue(); + AscendC::LocalTensor yLocal = inQueueY.DeQue(); + AscendC::LocalTensor zLocal = outQueueZ.AllocTensor(); + AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH); + outQueueZ.EnQue(zLocal); + inQueueX.FreeTensor(xLocal); + inQueueY.FreeTensor(yLocal); + } + __aicore__ inline void CopyOut(int32_t progress) + { + AscendC::LocalTensor zLocal = outQueueZ.DeQue(); + AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); + outQueueZ.FreeTensor(zLocal); + } + +private: + AscendC::TPipe pipe; + AscendC::TQue inQueueX, inQueueY; + AscendC::TQue outQueueZ; + AscendC::GlobalTensor xGm; + AscendC::GlobalTensor yGm; + AscendC::GlobalTensor zGm; +}; + +template +__global__ __aicore__ void add_custom(BOOL_1 bool_1, bool bool_2, uint32_t uint32_1, A tmpStruct, UINT16_T_1 uint16_1, GM_ADDR x, GM_ADDR y, GM_ADDR z) +{ + KernelAdd op; + op.Init(x, y, z); + op.Process(); +} + +#ifndef ASCENDC_CPU_DEBUG +void add_custom_do(uint32_t blockDim, void *stream, bool bool_1, bool bool_2, uint32_t uint32_1, A tmpStruct, uint16_t uint16_1, uint8_t *x, uint8_t *y, uint8_t *z) +{ + add_custom<<>>(bool_1, bool_2, uint32_1, tmpStruct, uint16_1, x, y, z); +} +#endif diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom_struct.h b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom_struct.h new file mode 100644 index 000000000..498e197cd --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/add_custom_struct.h @@ -0,0 +1,22 @@ +/** + * @file add_custom_tiling.h + * + * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. + */ +#ifndef ADD_CUSTOM_STRUCT_H +#define ADD_CUSTOM_STRUCT_H +#include + +struct A { + uint16_t A_uint16_1; + uint16_t A_uint64_1; + bool A_bool_1; + bool A_bool_2; + uint16_t A_uint16_2; +}; + +#endif // ADD_CUSTOM_STRUCT_H \ No newline at end of file diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/data_utils.h b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/data_utils.h new file mode 100644 index 000000000..09d906371 --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/data_utils.h @@ -0,0 +1,203 @@ +/** + * @file data_utils.h + * + * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. + */ +#ifndef DATA_UTILS_H +#define DATA_UTILS_H +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "acl/acl.h" + +typedef enum { + DT_UNDEFINED = -1, + FLOAT = 0, + HALF = 1, + INT8_T = 2, + INT32_T = 3, + UINT8_T = 4, + INT16_T = 6, + UINT16_T = 7, + UINT32_T = 8, + INT64_T = 9, + UINT64_T = 10, + DOUBLE = 11, + BOOL = 12, + STRING = 13, + COMPLEX64 = 16, + COMPLEX128 = 17, + BF16 = 27 +} printDataType; + +#define INFO_LOG(fmt, args...) fprintf(stdout, "[INFO] " fmt "\n", ##args) +#define WARN_LOG(fmt, args...) fprintf(stdout, "[WARN] " fmt "\n", ##args) +#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args) +#define CHECK_ACL(x) \ + do { \ + aclError __ret = x; \ + if (__ret != ACL_ERROR_NONE) { \ + std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \ + } \ + } while (0); + +/** + * @brief Read data from file + * @param [in] filePath: file path + * @param [out] fileSize: file size + * @return read result + */ +bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize) +{ + struct stat sBuf; + int fileStatus = stat(filePath.data(), &sBuf); + if (fileStatus == -1) { + ERROR_LOG("failed to get file"); + return false; + } + if (S_ISREG(sBuf.st_mode) == 0) { + ERROR_LOG("%s is not a file, please enter a file", filePath.c_str()); + return false; + } + + std::ifstream file; + file.open(filePath, std::ios::binary); + if (!file.is_open()) { + ERROR_LOG("Open file failed. path = %s", filePath.c_str()); + return false; + } + + std::filebuf *buf = file.rdbuf(); + size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in); + if (size == 0) { + ERROR_LOG("file size is 0"); + file.close(); + return false; + } + if (size > bufferSize) { + ERROR_LOG("file size is larger than buffer size"); + file.close(); + return false; + } + buf->pubseekpos(0, std::ios::in); + buf->sgetn(static_cast(buffer), size); + fileSize = size; + file.close(); + return true; +} + +/** + * @brief Write data to file + * @param [in] filePath: file path + * @param [in] buffer: data to write to file + * @param [in] size: size to write + * @return write result + */ +bool WriteFile(const std::string &filePath, const void *buffer, size_t size) +{ + if (buffer == nullptr) { + ERROR_LOG("Write file failed. buffer is nullptr"); + return false; + } + + int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE); + if (fd < 0) { + ERROR_LOG("Open file failed. path = %s", filePath.c_str()); + return false; + } + + size_t writeSize = write(fd, buffer, size); + (void)close(fd); + if (writeSize != size) { + ERROR_LOG("Write file Failed."); + return false; + } + + return true; +} + +template void DoPrintData(const T *data, size_t count, size_t elementsPerRow) +{ + assert(elementsPerRow != 0); + for (size_t i = 0; i < count; ++i) { + std::cout << std::setw(10) << data[i]; + if (i % elementsPerRow == elementsPerRow - 1) { + std::cout << std::endl; + } + } +} + +void DoPrintHalfData(const aclFloat16 *data, size_t count, size_t elementsPerRow) +{ + assert(elementsPerRow != 0); + for (size_t i = 0; i < count; ++i) { + std::cout << std::setw(10) << std::setprecision(6) << aclFloat16ToFloat(data[i]); + if (i % elementsPerRow == elementsPerRow - 1) { + std::cout << std::endl; + } + } +} + +void PrintData(const void *data, size_t count, printDataType dataType, size_t elementsPerRow = 16) +{ + if (data == nullptr) { + ERROR_LOG("Print data failed. data is nullptr"); + return; + } + + switch (dataType) { + case BOOL: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case INT8_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case UINT8_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case INT16_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case UINT16_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case INT32_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case UINT32_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case INT64_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case UINT64_T: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case HALF: + DoPrintHalfData(reinterpret_cast(data), count, elementsPerRow); + break; + case FLOAT: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + case DOUBLE: + DoPrintData(reinterpret_cast(data), count, elementsPerRow); + break; + default: + ERROR_LOG("Unsupported type: %d", dataType); + } + std::cout << std::endl; +} +#endif // DATA_UTILS_H diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/main.cpp b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/main.cpp new file mode 100644 index 000000000..dade08b62 --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/main.cpp @@ -0,0 +1,73 @@ +/** + * @file main.cpp + * + * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. + */ +#include "add_custom_struct.h" +#include "data_utils.h" +#include "acl/acl.h" +#include "aclrtlaunch_add_custom.h" + +int32_t main(int32_t argc, char *argv[]) +{ + uint32_t blockDim = 8; + size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); + size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); + + bool bool_1 = 1; + bool bool_2 = 1; + uint32_t uint32_1 = 37; + + A tmpStruct; + tmpStruct.A_uint16_1 = 23; + tmpStruct.A_uint64_1 = 43; + tmpStruct.A_bool_1 = true; + tmpStruct.A_bool_2 = true; + tmpStruct.A_uint16_2 = 277; + + uint16_t uint16_1 = 763; + + CHECK_ACL(aclInit(nullptr)); + int32_t deviceId = 0; + CHECK_ACL(aclrtSetDevice(deviceId)); + aclrtStream stream = nullptr; + CHECK_ACL(aclrtCreateStream(&stream)); + + uint8_t *xHost, *yHost, *zHost; + uint8_t *xDevice, *yDevice, *zDevice; + + CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize)); + CHECK_ACL(aclrtMallocHost((void **)(&yHost), inputByteSize)); + CHECK_ACL(aclrtMallocHost((void **)(&zHost), outputByteSize)); + CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); + CHECK_ACL(aclrtMalloc((void **)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); + CHECK_ACL(aclrtMalloc((void **)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST)); + + ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize); + ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize); + + CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); + CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE)); + + ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, bool_1, bool_2, uint32_1, tmpStruct, uint16_1, xDevice, yDevice, zDevice); + CHECK_ACL(aclrtSynchronizeStream(stream)); + + CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST)); + WriteFile("./output/output_z.bin", zHost, outputByteSize); + + CHECK_ACL(aclrtFree(xDevice)); + CHECK_ACL(aclrtFree(yDevice)); + CHECK_ACL(aclrtFree(zDevice)); + CHECK_ACL(aclrtFreeHost(xHost)); + CHECK_ACL(aclrtFreeHost(yHost)); + CHECK_ACL(aclrtFreeHost(zHost)); + + CHECK_ACL(aclrtDestroyStream(stream)); + CHECK_ACL(aclrtResetDevice(deviceId)); + CHECK_ACL(aclFinalize()); + return 0; +} diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/run.sh b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/run.sh new file mode 100755 index 000000000..9eb5de466 --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/run.sh @@ -0,0 +1,120 @@ +#!/bin/bash +CURRENT_DIR=$( + cd $(dirname ${BASH_SOURCE:-$0}) + pwd +) + +BUILD_TYPE="Debug" +INSTALL_PREFIX="${CURRENT_DIR}/out" + +SHORT=r:,v:,i:,b:,p:, +LONG=run-mode:,soc-version:,install-path:,build-type:,install-prefix:, +OPTS=$(getopt -a --options $SHORT --longoptions $LONG -- "$@") +eval set -- "$OPTS" +SOC_VERSION="Ascend310P3" + +while :; do + case "$1" in + -r | --run-mode) + RUN_MODE="$2" + shift 2 + ;; + -v | --soc-version) + SOC_VERSION="$2" + shift 2 + ;; + -i | --install-path) + ASCEND_INSTALL_PATH="$2" + shift 2 + ;; + -b | --build-type) + BUILD_TYPE="$2" + shift 2 + ;; + -p | --install-prefix) + INSTALL_PREFIX="$2" + shift 2 + ;; + --) + shift + break + ;; + *) + echo "[ERROR]: Unexpected option: $1" + break + ;; + esac +done + +RUN_MODE_LIST="cpu sim npu" +if [[ " $RUN_MODE_LIST " != *" $RUN_MODE "* ]]; then + echo "[ERROR]: RUN_MODE error, This sample only support specify cpu, sim or npu!" + exit -1 +fi + +VERSION_LIST="Ascend910A Ascend910B Ascend310B1 Ascend310B2 Ascend310B3 Ascend310B4 Ascend310P1 Ascend310P3 Ascend910B1 Ascend910B2 Ascend910B3 Ascend910B4" +if [[ " $VERSION_LIST " != *" $SOC_VERSION "* ]]; then + echo "[ERROR]: SOC_VERSION should be in [$VERSION_LIST]" + exit -1 +fi + +if [ -n "$ASCEND_INSTALL_PATH" ]; then + _ASCEND_INSTALL_PATH=$ASCEND_INSTALL_PATH +elif [ -n "$ASCEND_HOME_PATH" ]; then + _ASCEND_INSTALL_PATH=$ASCEND_HOME_PATH +else + if [ -d "$HOME/Ascend/ascend-toolkit/latest" ]; then + _ASCEND_INSTALL_PATH=$HOME/Ascend/ascend-toolkit/latest + else + _ASCEND_INSTALL_PATH=/usr/local/Ascend/ascend-toolkit/latest + fi +fi + +export ASCEND_TOOLKIT_HOME=${_ASCEND_INSTALL_PATH} +export ASCEND_HOME_PATH=${_ASCEND_INSTALL_PATH} +echo "[INFO]: Current compile soc version is ${SOC_VERSION}" +source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash +if [ "${RUN_MODE}" = "sim" ]; then + # in case of running op in simulator, use stub .so instead + export LD_LIBRARY_PATH=${_ASCEND_INSTALL_PATH}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH +elif [ "${RUN_MODE}" = "cpu" ]; then + export LD_LIBRARY_PATH=${_ASCEND_INSTALL_PATH}/tools/tikicpulib/lib:${_ASCEND_INSTALL_PATH}/tools/tikicpulib/lib/${SOC_VERSION}:${_ASCEND_INSTALL_PATH}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH +fi + +set -e +rm -rf build out +mkdir -p build +cmake -B build \ + -DRUN_MODE=${RUN_MODE} \ + -DSOC_VERSION=${SOC_VERSION} \ + -DCMAKE_BUILD_TYPE=${BUILD_TYPE} \ + -DCMAKE_INSTALL_PREFIX=${INSTALL_PREFIX} \ + -DASCEND_CANN_PACKAGE_PATH=${_ASCEND_INSTALL_PATH} +cmake --build build -j +cmake --install build + +rm -f ascendc_kernels_bbit +cp ./out/bin/ascendc_kernels_bbit ./ +rm -rf input output +mkdir -p input output +python3 scripts/gen_data.py +( + export LD_LIBRARY_PATH=$(pwd)/out/lib:$(pwd)/out/lib64:${_ASCEND_INSTALL_PATH}/lib64:$LD_LIBRARY_PATH + if [[ "$RUN_WITH_TOOLCHAIN" -eq 1 ]]; then + if [ "${RUN_MODE}" = "npu" ]; then + msprof op --application=./ascendc_kernels_bbit + elif [ "${RUN_MODE}" = "sim" ]; then + msprof op simulator --application=./ascendc_kernels_bbit + elif [ "${RUN_MODE}" = "cpu" ]; then + ./ascendc_kernels_bbit + fi + else + ./ascendc_kernels_bbit + fi +) +# tidy folder by delete log files +if [ "${RUN_MODE}" = "sim" ]; then + rm -f *.log *.dump *.vcd *.toml *_log +fi +md5sum output/*.bin +python3 scripts/verify_result.py output/output_z.bin output/golden.bin diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/gen_data.py b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/gen_data.py new file mode 100644 index 000000000..ea8ce828a --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/gen_data.py @@ -0,0 +1,25 @@ +#!/usr/bin/python3 +# coding=utf-8 +# +# Copyright (C) 2023-2024. Huawei Technologies Co., Ltd. All rights reserved. +# +# This program is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. +# =============================================================================== + +import numpy as np + + +def gen_golden_data_simple(): + input_x = np.random.uniform(1, 100, [8, 2048]).astype(np.float16) + input_y = np.random.uniform(1, 100, [8, 2048]).astype(np.float16) + golden = (input_x + input_y).astype(np.float16) + + input_x.tofile("./input/input_x.bin") + input_y.tofile("./input/input_y.bin") + golden.tofile("./output/golden.bin") + + +if __name__ == "__main__": + gen_golden_data_simple() diff --git a/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/verify_result.py b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/verify_result.py new file mode 100644 index 000000000..2dd46f803 --- /dev/null +++ b/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationTemplateNeo/scripts/verify_result.py @@ -0,0 +1,53 @@ +#!/usr/bin/python3 +# coding=utf-8 +# +# Copyright (C) 2023-2024. Huawei Technologies Co., Ltd. All rights reserved. +# +# This program is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. +# =============================================================================== + +import sys +import numpy as np + +# for float16 +relative_tol = 1e-3 +absolute_tol = 1e-5 +error_tol = 1e-3 + + +def verify_result(output, golden): + output = np.fromfile(output, dtype=np.float16).reshape(-1) + golden = np.fromfile(golden, dtype=np.float16).reshape(-1) + different_element_results = np.isclose(output, + golden, + rtol=relative_tol, + atol=absolute_tol, + equal_nan=True) + different_element_indexes = np.where(different_element_results == False)[0] + for index in range(len(different_element_indexes)): + real_index = different_element_indexes[index] + golden_data = golden[real_index] + output_data = output[real_index] + print( + "data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" % + (real_index, golden_data, output_data, + abs(output_data - golden_data) / golden_data)) + if index == 100: + break + error_ratio = float(different_element_indexes.size) / golden.size + print("error ratio: %.4f, tolerance: %.4f" % (error_ratio, error_tol)) + return error_ratio <= error_tol + + +if __name__ == '__main__': + try: + res = verify_result(sys.argv[1], sys.argv[2]) + if not res: + raise ValueError("[ERROR] result error") + else: + print("test pass") + except Exception as e: + print(e) + sys.exit(1) -- Gitee From 38096a1b90333a45c394c8a538c21b88b97dfbb1 Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Mon, 24 Nov 2025 17:07:30 +0800 Subject: [PATCH 8/9] update unaligned abs case --- .../CMakeLists.txt | 28 ++++++++++++------- .../cmake/cpu_lib.cmake | 9 ------ .../cmake/npu_lib.cmake | 11 -------- .../AbsDuplicateKernelInvocation/main.cpp | 19 ------------- .../CMakeLists.txt | 22 ++++++++++++--- .../cmake/cpu_lib.cmake | 9 ------ .../cmake/npu_lib.cmake | 11 -------- .../AbsGatherMaskKernelInvocation/main.cpp | 16 ----------- .../AbsPadKernelInvocation/CMakeLists.txt | 27 ++++++++++++------ .../cmake/cpu_lib.cmake | 9 ------ .../cmake/npu_lib.cmake | 11 -------- .../AbsPadKernelInvocation/main.cpp | 22 --------------- .../AbsUnPadKernelInvocation/CMakeLists.txt | 27 ++++++++++++------ .../cmake/cpu_lib.cmake | 10 ------- .../cmake/npu_lib.cmake | 11 -------- .../AbsUnPadKernelInvocation/main.cpp | 24 ---------------- 16 files changed, 72 insertions(+), 194 deletions(-) delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/npu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/npu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/npu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/CMakeLists.txt b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/CMakeLists.txt index 5216dbccc..4ca898d47 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/CMakeLists.txt +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/abs_duplicate.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_compile_options(ascendc_kernels_bbit PRIVATE @@ -33,14 +48,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} - graph_base - tiling_api - register - platform - ascendalog - dl + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/cpu_lib.cmake deleted file mode 100644 index 95a4d07f7..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/cpu_lib.cmake +++ /dev/null @@ -1,9 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) \ No newline at end of file diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/main.cpp b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/main.cpp index 3bd34e7a8..bc37be332 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/main.cpp +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsDuplicateKernelInvocation/main.cpp @@ -1,12 +1,7 @@ #include "data_utils.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_abs_duplicate_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void abs_duplicate_custom(GM_ADDR inputGM, GM_ADDR outputGM, GM_ADDR syncGM); -#endif int32_t main(int32_t argc, char *argv[]) { uint32_t blockDim = 4; @@ -18,19 +13,6 @@ int32_t main(int32_t argc, char *argv[]) uint32_t defaultSyncByteSizeSinleCore = 32; // 32B for SyncAll size_t syncByteSize = blockDim * defaultSyncByteSizeSinleCore; -#ifdef ASCENDC_CPU_DEBUG - uint8_t *inputGM = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *outputGM = (uint8_t *)AscendC::GmAlloc(outputByteSize); - uint8_t *syncGM = (uint8_t *)AscendC::GmAlloc(syncByteSize); - ReadFile("./input/input_x.bin", inputByteSize, inputGM, inputByteSize); - ReadFile("./input/sync.bin", syncByteSize, syncGM, syncByteSize); - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(abs_duplicate_custom, blockDim, inputGM, outputGM, syncGM); // use this macro for cpu debug - WriteFile("./output/output_z.bin", outputGM, outputByteSize); - AscendC::GmFree((void *)inputGM); - AscendC::GmFree((void *)outputGM); - AscendC::GmFree((void *)syncGM); -#else CHECK_ACL(aclInit(nullptr)); aclrtContext context; int32_t deviceId = 0; @@ -72,6 +54,5 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyContext(context)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif return 0; } diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/CMakeLists.txt b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/CMakeLists.txt index 9aed08955..466d94e11 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/CMakeLists.txt +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/abs_gather_mask.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_compile_options(ascendc_kernels_bbit PRIVATE @@ -33,8 +48,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/cpu_lib.cmake deleted file mode 100644 index 5362c8b5a..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/cpu_lib.cmake +++ /dev/null @@ -1,9 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/main.cpp b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/main.cpp index b5e01e4d0..c085687d0 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/main.cpp +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsGatherMaskKernelInvocation/main.cpp @@ -8,30 +8,15 @@ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. */ #include "data_utils.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_abs_gather_mask_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void abs_gather_mask_custom(GM_ADDR inputGM, GM_ADDR outputGM); -#endif int32_t main(int32_t argc, char *argv[]) { uint32_t blockDim = 8; size_t inputByteSize = 2318 * sizeof(int16_t); // 2318 = 2304 + 32 - 18 size_t outputByteSize = 2304 * sizeof(int16_t); -#ifdef ASCENDC_CPU_DEBUG - uint8_t *inputGM = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *outputGM = (uint8_t *)AscendC::GmAlloc(outputByteSize); - ReadFile("./input/input_x.bin", inputByteSize, inputGM, inputByteSize); - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(abs_gather_mask_custom, blockDim, inputGM, outputGM); // use this macro for cpu debug - WriteFile("./output/output_z.bin", outputGM, outputByteSize); - AscendC::GmFree((void *)inputGM); - AscendC::GmFree((void *)outputGM); -#else CHECK_ACL(aclInit(nullptr)); aclrtContext context; int32_t deviceId = 0; @@ -67,6 +52,5 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyContext(context)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif return 0; } diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/CMakeLists.txt b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/CMakeLists.txt index 783b44807..42d4bcbba 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/CMakeLists.txt +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/abs_pad.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp ${CMAKE_CURRENT_SOURCE_DIR}/abs_pad_tiling.cpp @@ -36,14 +51,8 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} + ascendc_kernels graph_base - tiling_api - register - platform - ascendalog - dl ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/cpu_lib.cmake deleted file mode 100644 index 95a4d07f7..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/cpu_lib.cmake +++ /dev/null @@ -1,9 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17 -D_GLIBCXX_USE_CXX11_ABI=0) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) \ No newline at end of file diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/main.cpp b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/main.cpp index 0582a8703..72d7a4338 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/main.cpp +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsPadKernelInvocation/main.cpp @@ -1,13 +1,8 @@ #include "data_utils.h" #include "kernel_tiling/kernel_tiling.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_abs_pad_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void abs_pad_custom(GM_ADDR inputGM, GM_ADDR outputGM, GM_ADDR tilingData); -#endif extern void GenerateTiling(const std::vector shapePad, const std::vector shapeUsed, uint8_t *tilingBuf); int32_t main(int32_t argc, char *argv[]) @@ -28,22 +23,6 @@ int32_t main(int32_t argc, char *argv[]) uint8_t *tilingBuf = (uint8_t *)malloc(tilingSize); GenerateTiling(shapePad, shapeUsed, tilingBuf); -#ifdef ASCENDC_CPU_DEBUG - uint8_t *inputGM = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *outputGM = (uint8_t *)AscendC::GmAlloc(outputByteSize); - uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tilingSize); - memcpy_s(tiling, tilingSize, tilingBuf, tilingSize); - ReadFile("./input/input_x.bin", inputByteSize, inputGM, inputByteSize); - - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(abs_pad_custom, blockDim, inputGM, outputGM, tiling); // use this macro for cpu debug - - WriteFile("./output/output_z.bin", outputGM, outputFileSize); - - AscendC::GmFree((void *)inputGM); - AscendC::GmFree((void *)outputGM); - AscendC::GmFree((void *)tiling); -#else CHECK_ACL(aclInit(nullptr)); aclrtContext context; int32_t deviceId = 0; @@ -84,7 +63,6 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyContext(context)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif free(tilingBuf); return 0; } diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/CMakeLists.txt b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/CMakeLists.txt index fc180c566..02adb55d8 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/CMakeLists.txt +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/abs_unpad.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp ${CMAKE_CURRENT_SOURCE_DIR}/abs_unpad_tiling.cpp @@ -40,14 +55,8 @@ target_compile_definitions(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} - tiling_api + ascendc_kernels graph_base - register - platform - ascendalog - dl ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/cpu_lib.cmake deleted file mode 100644 index feb68641b..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/cpu_lib.cmake +++ /dev/null @@ -1,10 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) - -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) diff --git a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/main.cpp b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/main.cpp index 93b0ea7da..7ef412af9 100644 --- a/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/main.cpp +++ b/operator/ascendc/0_introduction/16_unaligned_abs_kernellaunch/AbsUnPadKernelInvocation/main.cpp @@ -2,13 +2,8 @@ #include "kernel_tiling/kernel_tiling.h" #include "tiling/platform/platform_ascendc.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_abs_unpad_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void abs_unpad_custom(GM_ADDR inputGM, GM_ADDR outputGM, GM_ADDR tiling); -#endif extern void GenerateTiling(const std::vector shape, const char *socVersion, uint8_t *tilingBuf); int32_t main(int32_t argc, char *argv[]) @@ -29,24 +24,6 @@ int32_t main(int32_t argc, char *argv[]) uint8_t *tilingBuf = (uint8_t *)malloc(tilingSize); GenerateTiling(shape, socVersion, tilingBuf); -#ifdef ASCENDC_CPU_DEBUG - uint8_t *inputGM = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *outputGM = (uint8_t *)AscendC::GmAlloc(outputByteSize); - uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(sizeof(UnPadTiling)); - - ReadFile("./input/input_x.bin", inputByteSize, inputGM, inputByteSize); - memcpy_s(tiling, tilingSize, tilingBuf, tilingSize); - - AscendC::SetKernelMode(KernelMode::AIV_MODE); - - ICPU_RUN_KF(abs_unpad_custom, blockDim, inputGM, outputGM, tiling); // use this macro for cpu debug - - WriteFile("./output/output_z.bin", outputGM, outputByteSize); - - AscendC::GmFree((void *)inputGM); - AscendC::GmFree((void *)outputGM); - AscendC::GmFree((void *)tiling); -#else CHECK_ACL(aclInit(nullptr)); aclrtContext context; int32_t deviceId = 0; @@ -87,7 +64,6 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyContext(context)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif free(tilingBuf); return 0; } -- Gitee From 408f8725291ee0306e1d4f4e939ee472e8a90b51 Mon Sep 17 00:00:00 2001 From: zhangyujia77 Date: Mon, 24 Nov 2025 17:12:37 +0800 Subject: [PATCH 9/9] update unaligned reducemin case --- .../ReduceMinKernelInvocation/CMakeLists.txt | 22 +++++++++++++++---- .../cmake/cpu_lib.cmake | 9 -------- .../cmake/npu_lib.cmake | 11 ---------- .../ReduceMinKernelInvocation/main.cpp | 21 ------------------ 4 files changed, 18 insertions(+), 45 deletions(-) delete mode 100644 operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/cpu_lib.cmake delete mode 100644 operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/npu_lib.cmake diff --git a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/CMakeLists.txt b/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/CMakeLists.txt index cdcb350e3..61f089f1a 100644 --- a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/CMakeLists.txt +++ b/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/CMakeLists.txt @@ -18,12 +18,27 @@ endif() file(GLOB KERNEL_FILES ${CMAKE_CURRENT_SOURCE_DIR}/reduce_min.cpp) if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) + if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) + set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) + endif() + find_package(tikicpulib REQUIRED) elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) else() message("invalid RUN_MODE: ${RUN_MODE}") endif() + +if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) +elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) + set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) +else() + message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") +endif() +include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) + +# ascendc_library use to add kernel file to generate ascendc library +ascendc_library(ascendc_kernels SHARED ${KERNEL_FILES}) + add_executable(ascendc_kernels_bbit ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_compile_options(ascendc_kernels_bbit PRIVATE @@ -33,8 +48,7 @@ target_compile_options(ascendc_kernels_bbit PRIVATE target_link_libraries(ascendc_kernels_bbit PRIVATE $,$>:host_intf_pub>> - $:ascendcl>> - ascendc_kernels_${RUN_MODE} + ascendc_kernels ) install(TARGETS ascendc_kernels_bbit diff --git a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/cpu_lib.cmake b/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/cpu_lib.cmake deleted file mode 100644 index 5362c8b5a..000000000 --- a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/cpu_lib.cmake +++ /dev/null @@ -1,9 +0,0 @@ -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) -target_link_libraries(ascendc_kernels_${RUN_MODE} PUBLIC tikicpulib::${SOC_VERSION}) -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE -g -O0 -std=c++17) -install(TARGETS ascendc_kernels_${RUN_MODE} DESTINATION ${CMAKE_INSTALL_LIBDIR}) diff --git a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/npu_lib.cmake b/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/npu_lib.cmake deleted file mode 100644 index f92b095d1..000000000 --- a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/cmake/npu_lib.cmake +++ /dev/null @@ -1,11 +0,0 @@ -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -# ascendc_library use to add kernel file to generate ascendc library -ascendc_library(ascendc_kernels_${RUN_MODE} SHARED ${KERNEL_FILES}) diff --git a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/main.cpp b/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/main.cpp index 08c8d9ab0..fd23a84b8 100644 --- a/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/main.cpp +++ b/operator/ascendc/0_introduction/17_unaligned_reducemin_kernellaunch/ReduceMinKernelInvocation/main.cpp @@ -1,12 +1,7 @@ #include "data_utils.h" -#ifndef ASCENDC_CPU_DEBUG #include "acl/acl.h" #include "aclrtlaunch_reduce_min_custom.h" -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void reduce_min_custom(GM_ADDR inputGM, GM_ADDR outputGM, GM_ADDR syncGM); -#endif int32_t main(int32_t argc, char *argv[]) { uint32_t blockDim = 4; @@ -19,21 +14,6 @@ int32_t main(int32_t argc, char *argv[]) uint32_t defaultSyncByteSizeSinleCore = 32; // 32B for SyncAll size_t syncByteSize = blockDim * defaultSyncByteSizeSinleCore; -#ifdef ASCENDC_CPU_DEBUG - uint8_t *inputGM = (uint8_t *)AscendC::GmAlloc(inputByteSize); - uint8_t *outputGM = (uint8_t *)AscendC::GmAlloc(outputByteSize); - uint8_t *syncGM = (uint8_t *)AscendC::GmAlloc(syncByteSize); - ReadFile("./input/input_x.bin", inputByteSize, inputGM, inputByteSize); - ReadFile("./input/sync.bin", syncByteSize, syncGM, syncByteSize); - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(reduce_min_custom, blockDim, inputGM, outputGM, syncGM); // use this macro for cpu debug - - WriteFile("./output/output_z.bin", outputGM, outputByteSize); - - AscendC::GmFree((void *)inputGM); - AscendC::GmFree((void *)outputGM); - AscendC::GmFree((void *)syncGM); -#else CHECK_ACL(aclInit(nullptr)); aclrtContext context; int32_t deviceId = 0; @@ -75,6 +55,5 @@ int32_t main(int32_t argc, char *argv[]) CHECK_ACL(aclrtDestroyContext(context)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); -#endif return 0; } -- Gitee