{算子名} Kernel Launch 验证包【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime概述使用 内核调用符方式验证 {算子名} 自定义 Kernel 的 Runtime 调用流程。运算公式{写出算子的数学公式如 dst srcA alpha * srcB}目录说明{算子名小写}_verify/ ├── CMakeLists.txt find_package(ASC) 编译配置 ├── {算子名小写}.asc 单文件kernel host main ├── run.sh 一键编译脚本 └── README.md 本文件支持的产品型号Atlas A2 训练/推理系列产品默认 --npu-archdav-2201Atlas A3 训练/推理系列产品需改 --npu-archdav-3510前置条件Atlas A2 或 A3 系列 NPU 硬件CANN Toolkit 已安装含 ASC 编译器GCC 7.3, CMake 3.16环境变量设置# ${install_root} 替换为 CANN 安装根目录 source ${install_root}/cann/set_env.sh export ASCEND_INSTALL_PATH${install_root}/cann注意ASC 构建系统不需要 SOC_VERSION 和 ASCENDC_CMAKE_DIR 环境变量。修改目标 NPU 架构编辑 CMakeLists.txt注释/取消注释对应的 --npu-arch 行。快速使用bash run.sh ./build/main预期输出{根据核函数实际逻辑计算出的预期输出示例} Sample run successfully with kernel call!模板设计中的三个关键点值得注意 1. **运算公式必须写清**。精度验证的 expected 值正是依据该公式在 Host 侧计算的公式是验证逻辑正确性的基准线。 2. **环境变量大幅精简**。这是 ASC 构建系统相对 legacy 构建的核心优势只需 ASCEND_INSTALL_PATH 一个变量不再需要 SOC_VERSION 与 ASCENDC_CMAKE_DIRNPU 架构改由 CMakeLists.txt 中的 --npu-arch 指定见第 3 节。 3. **产品型号与架构映射**。Atlas A2 系列默认 dav-2201Atlas A3 系列使用 dav-3510通过注释/取消注释 CMakeLists.txt 中的编译选项行切换。 作为对照仓库中 legacy 方式ascendc_library()的官方示例 [0_launch_kernel/README.md](https://link.gitcode.com/i/af3c066dcfefb29b82edddb8ccacbe63) 需要先执行 source ${install_root}/cann/set_env.sh 再执行 source ${git_clone_path}/example/set_sample_env.sh 来自动识别 SOC_VERSION 和 ASCENDC_CMAKE_DIR其 [run.sh](https://link.gitcode.com/i/1fc75a53bcdd2a5b339c009fb85a5c5e) 中也存在对 SOC_VERSION、ASCENDC_CMAKE_DIR、ascendc.cmake 存在性的三重强制检查。验证包采用 ASC 构建系统后这些繁琐的环境配置全部消失这正是 ASC 构建系统开箱即用特性的直接体现。 --- ## 3. CMakeLists.txt 模板用 find_package(ASC) 替代 ascendc_library() ### 3.1 基础模板 cmake cmake_minimum_required(VERSION 3.16) find_package(ASC REQUIRED) project({算子名}_Verify LANGUAGES ASC CXX) add_executable(main {算子名小写}.asc) target_link_libraries(main PRIVATE m dl ) # NPU architecture: adjust to your hardware target_compile_options(main PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-archdav-2201 # Atlas A2 / Ascend 910B默认 # $$COMPILE_LANGUAGE:ASC:--npu-archdav-3510 # Atlas A3 / Ascend 950 )模板要点find_package(ASC REQUIRED)是 ASC 构建系统的入口它会自动完成编译器定位、库搜索路径设置等工作——这正是只需ASCEND_INSTALL_PATH一个环境变量的底层原因。project(... LANGUAGES ASC CXX)显式声明.asc语言使.asc单文件中的核函数部分由 ASC 编译器编译、Host 部分由 C 编译器编译。add_executable(main {算子名小写}.asc)直接将.asc文件作为可执行程序的唯一源文件——kernel 与 host 合一的单文件模式。--npu-arch通过target_compile_options配合$$COMPILE_LANGUAGE:ASC:...生成器表达式传入保证该选项只作用于 ASC 编译阶段不影响 Host C 编译。3.2 需要 tiling 等库时的扩展当算子涉及 tiling、Cube/Matmul 等高级特性时按需追加链接库target_link_libraries(main PRIVATE tiling_api register platform m dl )需要说明的是CANN 正确安装后ASC 构建系统find_package(ASC)会自动设置这些库的搜索路径开发者只需在 CMakeLists.txt 中声明依赖即可自然链接成功无需关心这些库的物理位置。3.3 与 legacy 构建方式的对照仓库中 0_launch_kernel/CMakeLists.txt 展示了 legacy 方式的典型写法set(ASCENDC_CMAKE_DIR $ENV{ASCENDC_CMAKE_DIR}) set(SOC_VERSION $ENV{SOC_VERSION}) include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_fatbin_library(ascendc_kernels_simple ${KERNEL_FILES_SIMPLE}) include_directories(${ASCEND_CANN_PACKAGE_PATH}/include) link_directories(${ASCEND_CANN_PACKAGE_PATH}/lib64)可见 legacy 方式需要手动设置ASCENDC_CMAKE_DIR、SOC_VERSION、ASCEND_CANN_PACKAGE_PATH三个环境变量且核函数与 Host 代码分离编译核函数用ascendc_fatbin_library编译为 fatbinHost 单独链接libacl_rt.so。验证包统一采用find_package(ASC)的现代方式原因有二配置更简单无需三个环境变量无需手动指定 include/link 路径避开ascendc_library()的已知缺陷该 legacy 宏的 auto_gen 机制会对核函数做宏重命名#define KernelName KernelName_origin导致__cube__核函数在 bisheng 编译器严格入口校验下报错即使是纯__aicore__核函数重命名也会在 ASC 编译器下触发auto type derivate failed。ASC 构建系统不存在此问题。禁止事项强制验证包中禁止使用ascendc_library()必须使用find_package(ASC).asc单文件模式。4..asc单文件模板真实核函数 Host main 的合体.asc文件是验证包的核心将核函数与 Host main 合并于同一文件。其模板与配套的强制规则如下。4.1 文件结构总览/** * {算子名} verification using ASC build system. * Single-file (.asc): REAL kernel host in one file. * Kernel code copied from reference/asc-devkit/, ReadFile/WriteFile replaced with hardcoded data. */ #include acl/acl.h #include kernel_operator.h #include data_utils.h // 按需添加其他头文件如 tiling/tiling_api.h, kernel_tiling/kernel_tiling.h 等 #include cstdint #include cstdio #include cstring #include vector // Kernel // 从 reference/asc-devkit/ 完整复制核函数代码 // 包括核函数类、辅助函数、tiling 生成函数等 // 保持原始的 __vector__/__cube__/__mix__(m,n) 属性标注 // Host #define CHECK_ERROR(ret) \ if ((ret) ! ACL_SUCCESS) { \ printf(Error at line %d, ret %d\n, __LINE__, ret); \ return -1; \ } int main() { const int32_t deviceId 0; const uint32_t blockDim /* 根据算子调整 */; aclrtStream stream nullptr; // 硬编码非零输入数据禁止全零否则无法验证精度 // std::vector... srcHost {1.0f, 2.0f, 3.0f, ...}; // std::vector... dstHost(..., 0); CHECK_ERROR(aclInit(nullptr)); printf(ACL init successfully\n); CHECK_ERROR(aclrtSetDevice(deviceId)); printf(Set device %d successfully\n, deviceId); CHECK_ERROR(aclrtCreateStream(stream)); printf(Create stream successfully\n); // 分配设备内存 uint8_t* srcDevice nullptr; uint8_t* dstDevice nullptr; // CHECK_ERROR(aclrtMalloc(...)); // 拷贝输入数据到设备 // CHECK_ERROR(aclrtMemcpy(..., ACL_MEMCPY_HOST_TO_DEVICE)); // 直接调用ASC 单文件模式不需要包装函数 {算子名}KernelblockDim, nullptr, stream(srcDevice, dstDevice); printf(Custom AscendC kernel call successfully\n); CHECK_ERROR(aclrtSynchronizeStream(stream)); printf(Synchronize stream successfully\n); // 结果回传 // CHECK_ERROR(aclrtMemcpy(..., ACL_MEMCPY_DEVICE_TO_HOST)); // 逐元素打印结果expected 必须与核函数实际逻辑一致 // printf( result[%u] %.1f (expected: %.1f)\n, i, dstHost[i], expected); // 释放资源 // CHECK_ERROR(aclrtFree(srcDevice)); // CHECK_ERROR(aclrtFree(dstDevice)); CHECK_ERROR(aclrtDestroyStream(stream)); printf(Destroy stream successfully\n); CHECK_ERROR(aclrtResetDeviceForce(deviceId)); printf(Reset device successfully\n); CHECK_ERROR(aclFinalize()); printf(ACL finalize successfully\n); printf(\nSample run successfully with kernel call!\n); return 0; }4.2 Host 侧调用链与仓库示例逐段印证模板中的 Host 流程与仓库示例 0_launch_kernel/main.cpp 中InitializeRuntimemain.cpp L59-L66的调用顺序完全一致aclInit(nullptr)→aclrtSetDevice(0)→aclrtCreateStream(stream)。资源释放侧模板使用aclrtDestroyStreamaclrtResetDeviceForceaclFinalize仓库示例则在ReleaseKernelResourcesmain.cpp L177-L212中以aclrtDestroyStreamForceaclrtResetDeviceForceaclFinalize顺序完成清理并配套UpdateFinalResultOnError逐 API 记录错误码。两者体现同一原则每个 ACL API 的返回值都必须检查模板中的CHECK_ERROR宏即为此设计。blockDim对应的并行度使用的 AI Core 数量需根据算子并行策略调整。仓库示例使用 8 核并行main.cpp L216核函数内部按 8 核切分数据可作为取值参考。4.3 核函数规则强制项核函数部分是验证包与玩具 demo的本质区别。以下是必须严格遵守的规则规则要求原因使用真实核函数从reference/asc-devkit/完整复制算子核函数代码含核函数类、辅助函数、tiling 生成函数禁止简化为 element-wise copy简化核函数跳过 tiling 参数准备、workspace 分配、多参数调用等真实环节断点识别不完整属性保持一致__vector__/__cube__/__mix__(m,n)与原始代码一致原始代码若用extern C则保留保证核函数类型正确推导与编译禁止__aicore__不得使用__aicore__属性ASC 下无法推导核函数类型报错auto type derivate failed唯一允许的修改ReadFile/WriteFile替换为硬编码数据初始化 printf输出保证验证包自包含、可独立运行辅助头文件一并复制data_utils.h、nd2nz_utils.h等随核函数复制到验证包保证编译不缺失头文件关于__aicore__的禁止原因参考 kernel-function-and-launch-syntax.md 的说明__aicore__是ascendc_library()的旧写法在 ASC 构建系统下ASC 编译器无法从__aicore__推导核函数类型——若函数体内无 AscendC API 调用会直接报kernel type of __global__ func not marked. auto type derivate may be failed。对比仓库 legacy 示例 add_custom.cpp 中extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)的写法可见验证包在 ASC 下应改用显式类型属性。4.4 核函数参数与调用语法是 AscendC 扩展的内核调用符用于从 Host 侧发起核函数执行kernel_nameblockDim, l2ctrl, stream(参数列表);三个控制参数的含义与取值如下参数类型含义取值范围blockDimuint32_t核函数使用的核数并行度[1, 65535]l2ctrlvoid*保留参数L2 缓存控制固定传nullptrstreamaclrtStream执行流任务队列由aclrtCreateStream创建调用是异步的立即返回核函数在 Device 侧异步执行必须通过aclrtSynchronizeStream(stream)等待完成。同一 Stream 上的多个核函数按提交顺序FIFO执行不同 Stream 上的核函数可并行。Device 指针类型匹配⚠️ 易错点ASC 编译器在调用时只做T*→__gm__ T*的地址空间转换不做跨类型转换。因此 Device 指针的 C 类型必须与核函数参数去掉__gm__后的类型严格一致核函数参数为GM_ADDR即__gm__ uint8_t*→ Device 指针必须声明为uint8_t*核函数参数为__gm__ float*→ Device 指针必须声明为float*禁止将float*传给GM_ADDR参数报错cannot initialize __gm__ uint8_t * with float *禁止reinterpret_castGM_ADDR(floatPtr)报错reinterpret_cast from float * to __gm__ uint8_t * is not allowed5. run.sh 模板一键编译脚本#!/bin/bash set -e if [ -z ${ASCEND_INSTALL_PATH} ]; then echo [ERROR]: ASCEND_INSTALL_PATH is not set. Please source \${install_root}/cann/set_env.sh first. exit 1 fi if [ ! -f ${ASCEND_INSTALL_PATH}/bin/setenv.bash ]; then echo [ERROR]: ${ASCEND_INSTALL_PATH}/bin/setenv.bash does not exist. exit 1 fi _ASCEND_INSTALL_PATH$ASCEND_INSTALL_PATH source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash SCRIPT_DIR$(cd $(dirname ${BASH_SOURCE[0]}) pwd) cd ${SCRIPT_DIR} BUILD_DIR${SCRIPT_DIR}/build mkdir -p ${BUILD_DIR} cd ${BUILD_DIR} echo Configuring CMake (ASC build system)... cmake .. echo Building... make -j$(nproc) cd ${SCRIPT_DIR} echo Build completed successfully! echo Executable location: ${SCRIPT_DIR}/build/main echo Run the sample: ./build/main脚本设计遵循最小化检查 最大化自动化原则仅检查ASCEND_INSTALL_PATH未设置则提示先 sourceset_env.sh检查setenv.bash存在性该文件是 CANN Toolkit 的环境初始化脚本位于${ASCEND_INSTALL_PATH}/bin/下存在与否是 CANN 安装是否完整的最直接信号source setenv.bash加载编译环境补充 ASC 编译器、CMake 工具链等的路径cmake ..make不需要-DASCEND_CANN_PACKAGE_PATH等额外参数——这也是 ASC 构建系统相对 legacy 方式的又一处简化。作为对照legacy 示例的 run.sh 在 cmake 阶段需要显式传入-DSOC_VERSION、-DASCEND_CANN_PACKAGE_PATH等选项。6. 精度验证让验证包真正验证正确性验证包区别于一般能跑通demo 的核心在于不仅验证 Runtime 调用流程走通还验证计算结果的正确性。精度验证逻辑应嵌入.asc文件的 Host main 中按以下流程实现。6.1 构造非零输入数据将原始代码中的ReadFile替换为硬编码数据初始化禁止全零输入全零会导致大多数算子输出全零无法验证精度。推荐模式递增序列1, 2, 3, ...、全 1、小随机整数等。6.2 Host 侧计算 expected 值根据算子的数学公式在 main 函数中用 C 计算 expected 输出矩阵乘法 → 三重循环实现向量加法 → 逐元素相加排序算子 →std::sort复杂算子如 Matmul LeakyReLU 融合、NZ 格式转换→ 优先从reference/的scripts/gen_data.py中提取计算逻辑完整翻译为 C包括格式转换、后处理不得跳过或简化。6.3 逐元素对比与容差判定将 Device 返回结果与 expected 逐元素对比采用相对误差 绝对误差双重判定数据类型rtol相对误差容差atol绝对误差容差float321e-31e-3float16/half1e-21e-2int精确匹配精确匹配判定规则|actual - expected| atol rtol * |expected|。打印前几个元素的 actual vs expected然后输出统计结果通过[PRECISION PASS] X/Y elements passed (max_diffZ) 失败[PRECISION FAIL] X/Y elements failed (max_diffZ)6.4 验证通过条件程序必须同时满足以下两个条件才算验证通过输出Sample run successfully with kernel call!精度验证通过输出[PRECISION PASS]且通过率 ≥ 99%若精度不达标即使程序正常运行也必须输出[PRECISION FAIL]并报告失败元素数和最大误差以便定位 expected 计算错误、数据类型不匹配或 tiling 参数问题。这种运行成功 精度达标双判定思路与仓库 legacy 示例的做法一脉相承0_launch_kernel 示例运行后通过python3 scripts/gen_data.py生成输入与 golden 数据执行 kernel 后用 verify_result.py 对比输出得到error ratio: 0.0000, tolerance: 0.0010与[SUCCESS] result correct的结果。验证包把这条外部脚本校验链路内化到了.asc单文件的 Host 代码中做到自包含、单文件、可拷贝即跑。7. 注意事项速查表构建验证包前的自检清单ASC 构建系统— 使用find_package(ASC).asc单文件禁止使用ascendc_library()。真实核函数强制— 从reference/asc-devkit/完整复制算子核函数禁止简化为 element-wise copy 或其他简化逻辑核函数属性与原始代码一致__vector__/__cube__/__mix__(m,n)原始代码有extern C则保留禁止__aicore__。精度验证强制— 验证包不仅验证 Runtime 调用流程还必须在 Host 侧计算 expected 值并与 Device 输出逐元素对比相对误差 绝对误差双重判定float32: rtol1e-3/atol1e-3half: rtol1e-2/atol1e-2int: 精确匹配输出[PRECISION PASS]或[PRECISION FAIL]。数据初始化— 将ReadFile替换为硬编码非零数据递增序列、全 1、小随机整数等禁止全零WriteFile替换为精度对比 printf输出这是唯一允许对原始.asc代码做的修改。辅助头文件— 如算子依赖data_utils.h、nd2nz_utils.h等必须从reference/一并复制到验证包中。环境变量— 仅需ASCEND_INSTALL_PATH不需要SOC_VERSION和ASCENDC_CMAKE_DIR。NPU 架构— 在 CMakeLists.txt 中通过--npu-arch指定Atlas A2 用dav-2201Atlas A3 用dav-3510不通过环境变量。Device 指针类型必须与核函数参数类型匹配— ASC 编译器在调用时只做地址空间转换添加__gm__不做类型转换GM_ADDR__gm__ uint8_t*参数必须用uint8_t*指针传入禁止跨类型传参或reinterpret_castGM_ADDR(...)。8. 生成后的强制验证编译与运行验证包文件全部写入后必须在真实环境中实际编译并运行不得以编译通过代替运行验证。执行流程如下# 1. 定位 CANN 安装路径ASCEND_INSTALL_PATH / 常见安装目录 export ASCEND_INSTALL_PATH${CANN_PATH} source ${CANN_PATH}/bin/setenv.bash # 2. 编译 cd reports/{算子名}_verify/ mkdir -p build cd build cmake .. make -j$(nproc) # 3. 运行 cd .. ./build/main【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考