CANN 昇腾算子直调实战Ascend C Add 算子样例CopyIn/Compute/CopyOut 三段式编程与双缓冲详解【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer本文基于开源仓库cann-recipes-harmony-infer中的 AddKernelInvocation 样例工程位于 ops/ascendc/AddKernelInvocation面向希望在鸿蒙昇腾如 KirinX90 / Kirin9030设备上通过 Ascend C 实现并直调自定义算子的开发者。文章将完整讲解样例的工程结构、Add 算子的 Kernel 实现含多核分块、双缓冲、TPipe/TQue 队列流水、CPU 调测与 NPU 直调两种调用方式以及把样例迁移替换为自有算子时需要修改的全部要点。读完本文你将掌握一条写 Kernel → 编译 → 生成数据 → 运行 → 真值比对的完整算子开发闭环。一、样例工程概览与目录结构AddKernelInvocation是一个算子直调不经过算子注册与推理框架直接在应用程序中调用 Kernel 函数的最小可运行样例固定实现 shape 为8 * 2048的 Add 算子z x y输入输出数据类型为halfFP16。它同时保留了 CPU 侧调测与 NPU 侧运行两条验证路径是理解 Ascend C 算子编程模型的最佳入门工程。其完整目录结构如下原样继承自 README.md├── AddKernelInvocation │ ├── cmake // 编译工程文件含 CCE 编译器适配与 npu 目标构建 │ ├── input // 存放脚本生成的输入数据目录 │ ├── output // 存放算子运行输出数据和真实数据golden的目录 │ ├── scripts │ │ ├── acl.json // acl配置文件本样例为空对象 {} │ │ ├── gen_data.py // 生成输入数据与真值golden数据的脚本 │ │ ├── verify_result.py // 对比算子输出与真值是否一致的验证脚本 │ │── add_custom.cpp // 算子kernel实现暴露给main.cpp调用 │ │── CMakeLists.txt // 编译工程文件 │ │── data_utils.h // 数据读入写出函数ReadFile / WriteFile / CHECK_ACL │ │── main.cpp // 主函数调用算子的应用程序含CPU域及NPU域调用 │ ├── run.sh // 编译运行算子的脚本从源码结构看各文件职责边界清晰add_custom.cpp只负责 Kernel 本身核函数与调用封装main.cpp负责数据准备、设备管理与内核发起data_utils.h提供与设备无关的二进制文件读写scripts/下两个 Python 脚本负责造数据与验结果run.sh将编译、运行、验证串成一条命令。二、Kernel 实现解析CopyIn / Compute / CopyOut 三段式编程2.1 数学表达式与片上执行模型Add 算子的数学表达式为z x y看似简单但它在昇腾设备上的执行模型与 CPU 完全不同Ascend C 提供的矢量计算接口如AscendC::Add操作的元素都是LocalTensor片上 Local Memory 中的张量而输入输出数据本身位于 Global Memory外部存储。因此正确流程是先把输入数据搬运进片上存储再调用计算接口完成加法最后把结果搬出到外部存储。2.2 多核切分与双缓冲参数设计Kernel 在 add_custom.cpp 顶部通过一组constexpr定义了数据规模与并行策略常量值含义TOTAL_LENGTH8 * 2048 16384数据总长度与8*2048的 shape 对应USE_CORE_NUM8使用的 AI Core 数量BLOCK_LENGTHTOTAL_LENGTH / USE_CORE_NUM 2048每个核负责的数据长度TILE_NUM8每个核的数据再切分为 8 个 tileBUFFER_NUM2每个队列的缓冲区个数双缓冲TILE_LENGTHBLOCK_LENGTH / TILE_NUM / BUFFER_NUM 128单次搬入/计算/搬出的数据长度其中TILE_LENGTH 128是由2048 / 8 / 2推导而来由于启用了BUFFER_NUM 2的双缓冲同一份 2048 长度数据被拆成 2 份交替使用每份再切 8 个 tile 流水最终每次处理 128 个元素。整个 Kernel 的循环次数为TILE_NUM * BUFFER_NUM 16次见 add_custom.cpp 的Process()每次循环依次执行 CopyIn → Compute → CopyOut形成搬运与计算重叠的流水线。2.3 KernelAdd 类的四个方法核心类KernelAdd定义于 add_custom.cpp四个方法分别对应1Init绑定全局内存与初始化队列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));通过AscendC::GetBlockIdx()获取当前核的编号将三个全局张量的起始地址各自偏移BLOCK_LENGTH * blockIdx实现多核数据切分——每个核只处理属于自己的 2048 个元素pipe.InitBuffer为输入队列inQueueX、inQueueY与输出队列outQueueZ各分配BUFFER_NUM2块、每块TILE_LENGTH * sizeof(half)字节的片上缓冲区即开启双缓冲。2CopyInGlobal Memory → Local MemoryAscendC::LocalTensorhalf xLocal inQueueX.AllocTensorhalf(); AscendC::LocalTensorhalf yLocal inQueueY.AllocTensorhalf(); AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal);从输入队列申请空闲缓冲区用AscendC::DataCopy把 Global Memory 上xGm、yGm的第progress * TILE_LENGTH段搬入再通过EnQue将填好的张量送入队列等待计算。3Compute执行矢量加法AscendC::LocalTensorhalf xLocal inQueueX.DeQuehalf(); AscendC::LocalTensorhalf yLocal inQueueY.DeQuehalf(); AscendC::LocalTensorhalf zLocal outQueueZ.AllocTensorhalf(); AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH); outQueueZ.EnQuehalf(zLocal); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal);从输入队列DeQue取出待计算张量调用矢量计算接口AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH)完成zLocal xLocal yLocal将结果EnQue到输出队列同时FreeTensor归还输入缓冲区以便下一轮 CopyIn 复用——这正是双缓冲流水得以循环运转的关键。4CopyOutLocal Memory → Global MemoryAscendC::LocalTensorhalf zLocal outQueueZ.DeQuehalf(); AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); outQueueZ.FreeTensor(zLocal);从输出队列取出计算结果用DataCopy写回 Global Memory 上的zGm最后释放输出缓冲区。2.4 核函数入口与对 main 暴露的接口Kernel 入口为extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)见 add_custom.cpp内部依次完成KernelAdd对象的Init与Process。在非 CPU 调测分支#ifndef ASCENDC_CPU_DEBUG下文件还封装了给main.cpp调用的普通 C 函数add_custom_doadd_custom.cppvoid add_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *z) { add_customblockDim, l2ctrl, stream(x, y, z); }这里使用昇腾的内核调用符blockDim, l2ctrl, stream发起 NPU 侧内核执行blockDim指定使用的核数样例为 8stream指定 ACL 流。通过这层包装main.cpp无需感知内核调用符语法只需声明extern void add_custom_do(...)即可完成直调。三、调用实现CPU 调测与 NPU 运行双路径应用程序 main.cpp 通过宏ASCENDC_CPU_DEBUG区分代码逻辑运行于 CPU 侧还是 NPU 侧两条路径在main()内由#ifdef编译期隔离3.1 CPU 侧运行验证ICPU_RUN_KF 调测宏#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); WriteFile(./output/output_z.bin, z, outputByteSize); AscendC::GmFree((void *)x); AscendC::GmFree((void *)y); AscendC::GmFree((void *)z);CPU 侧主要通过 CPU 调测库的ICPU_RUN_KF调测宏在 CPU 上模拟执行核函数先用AscendC::GmAlloc在模拟的 Global Memory 上分配内存ReadFile读入输入二进制SetKernelMode(KernelMode::AIV_MODE)指定矢量核AIV模式再以ICPU_RUN_KF(add_custom, blockDim, x, y, z)触发执行最后WriteFile落盘输出。需要说明的是本样例面向 Kirin 平台main.cpp 中明确注释Kirin暂不支持该编译分支且 run.sh 中亦有Kirin Soc Currently not support cpu!的校验拦截。因此在 KirinX90 / Kirin9030 上实际可用的运行方式是simNPU 仿真。3.2 NPU 侧运行验证ACL 直调链路NPU 分支main.cpp走标准 ACLAscend Computing Language调用链完整流程为aclInit(./scripts/acl.json)初始化 ACL 运行时配置文件本样例为空{}即全部使用默认配置aclrtSetDevice(deviceId)指定设备、aclrtCreateStream(stream)创建流aclrtMallocHost/aclrtMalloc分别申请主机侧与设备侧内存inputByteSize 8 * 2048 * sizeof(uint16_t) 32768字节half即 16 位ReadFile读入input_x.bin、input_y.bin经aclrtMemcpy以ACL_MEMCPY_HOST_TO_DEVICE拷贝到设备调用add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice)发起内核执行随后aclrtSynchronizeStream(stream)同步等待执行完成aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_HOST将结果拷回主机WriteFile写出output_z.bin依次释放设备/主机内存、销毁流、复位设备、aclFinalize()收尾。所有 ACL 调用均通过 data_utils.h 中的CHECK_ACL宏统一检查返回值非ACL_ERROR_NONE时打印出错文件与行号保证错误可定位。四、编译、运行与结果验证4.1 打开样例目录并配置环境变量cd $HOME/ops/AddKernelInvocation说明这里的$HOME需要替换为本仓库根目录即实际路径为ops/ascendc/AddKernelInvocation。export ASCEND_INSTALL_PATH$HOME/Ascend/ascend-toolkit/latest此外从 run.sh 的源码可以看到脚本对环境变量的探测顺序为优先读取ASCEND_TOOLKIT_HOME其次ASCEND_HOME_PATH再次默认路径$HOME/Ascend/ascend-toolkit/latest最后兜底/usr/local/Ascend/ascend-toolkit/latest也支持通过-i/--install-path参数直接传入。4.2 样例执行命令bash run.sh -r [RUN_MODE] -v [SOC_VERSION]两个核心参数说明同时可参考 run.sh 的 getopt 解析逻辑-r / --run-mode运行模式。本样例实际支持cpu / sim / npu三值见 run.sh 的RUN_MODE_LIST校验其中simNPU 仿真是 README 推荐的验证方式cpu模式在 Kirin 平台被显式禁用。-v / --soc-version芯片型号取值为KirinX90或Kirin9030脚本会通过VersionMap校验合法性并据此将CORE_TYPE置为AiCore。-i / --install-path可选直接指定 CANN 安装路径。典型示例bash run.sh -r sim -v KirinX904.3 run.sh 背后编译 → 造数 → 执行 → 验证以sim模式为例脚本实际做了五件事准备仿真环境将LD_LIBRARY_PATH指向$_ASCEND_INSTALL_PATH/x86_64-linux/simulator/${_SOC_VERSION}/lib并设置CAMODEL_LOG_PATH./sim_log见 run.shCMake 编译执行cmake -B build -Dsmoke_testcaseadd -DASCEND_PRODUCT_TYPEKirinX90 -DASCEND_CORE_TYPEAiCore -DASCEND_RUN_MODEsim -DASCEND_INSTALL_PATH...后cmake --build build --target add_sim产出可执行文件add_sim见 run.sh。CMake 侧由 cmake/npu/CMakeLists.txt 将*.cpp以 CCE 语言编译并链接libascendcl生成数据python3 scripts/gen_data.py执行算子运行./add_sim结果验证python3 scripts/verify_result.py output/output_z.bin output/golden.bin。4.4 数据生成与真值比对原理gen_data.py 使用 NumPy 生成 shape 为[8, 2048]、取值范围uniform(1, 100)的 FP16 随机数作为输入input_x、input_y并同步计算真值golden input_x input_y分别落盘到./input/input_x.bin、./input/input_y.bin与./output/golden.bininput_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)verify_result.py 将算子输出output_z.bin与真值golden.bin按 FP16 读入用np.isclose做逐元素近似比较rtol1e-3、atol1e-5容忍equal_nanTrue逐条打印前 100 个不一致元素的下标、期望值、实际值与相对偏差最后以不一致元素占比 ≤ 1e-3作为通过阈值输出test pass。五、替换为自定义算子四步迁移指南如果将该样例工程替换为自己的算子实现需要修改以下内容README 原文建议在run.sh、gen_data.py、main.cpp中搜索注释迁移算子修改点源码中对应的标记分别位于 run.sh、gen_data.py、main.cpp 与 main.cpp1. 替换add_custom.cpp整个文件该文件是核函数实现整个文件都要替换成新算子实现包括新的__global__ __aicore__核函数入口、KernelAdd或其等价类中的Init / CopyIn / Compute / CopyOut逻辑、以及在非 CPU 调测分支下向main.cpp暴露的xxx_do封装函数内部使用内核调用符。若新算子有多个输入、多个输出需同步调整Init中SetGlobalBuffer的参数个数与main侧接口签名。2. 修改run.sh的FILE_NAME在 run.sh 处根据实际算子名修改FILE_NAME属性默认add该变量控制编译获得的二进制名cmake --build build --target ${FILE_NAME}_${RUN_MODE}与后续执行文件./${FILE_NAME}_${RUN_MODE}均依赖它。3. 修改scripts/gen_data.py需自定义生成输入数据的内容包括输入数据的Shape当前为[8, 2048]、数据类型当前为float16需与 Kernel 内half及main.cpp中sizeof(uint16_t)保持一致、算子的输入参数数量当前为 x、y 两个输入并相应修改 golden 的计算公式。4. 修改main.cpp两处迁移算子修改点顶部extern void add_custom_do(...)声明改为调用add_custom.cpp中暴露的新符号见 main.cpp主体读数据与调用逻辑若直调的算子类型改变需结合算子实际输入和gen_data.py同步修改inputByteSize大小、读入文件的数量与顺序、add_custom_do的实参列表见 main.cppCPU 调测分支下同样要调整ReadFile的文件名与ICPU_RUN_KF的入参。完成上述修改后重新执行bash run.sh -r sim -v KirinX90即可完成新算子从编译到验证的全流程。六、小结与延伸阅读AddKernelInvocation样例完整展示了昇腾 Ascend C 算子开发的最小闭环多核切分 TPipe/TQue 队列 双缓冲流水CopyIn → Compute → CopyOut ACL 直调 Python 数据生成与真值比对。理解该样例后读者可以在本仓库 ops/ascendc/src 下找到多个基于相同三段式框架演进的更复杂算子实现如 sobel_custom、rmsNorm_custom、slice_gelu_custom、gather_dequant_int8_custom、quant_matmul_custom 等这些工程在op_hostTiling与op_kernelKernel分离的基础上进一步演示了算子在 ONNX 模型推理链路中的注册与使用方式可作为从直调走向框架集成的下一站参考对应的算子设计文档见 ops/ascendc/docs。若要在鸿蒙端侧完整跑通模型推理可进一步阅读仓库顶层 README.md 与 harmony_infer 目录下的端侧推理示例。【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN 昇腾算子直调实战：Ascend C Add 算子样例（CopyIn/Compute/CopyOut 三段式编程与双缓冲详解）
CANN 昇腾算子直调实战Ascend C Add 算子样例CopyIn/Compute/CopyOut 三段式编程与双缓冲详解【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer本文基于开源仓库cann-recipes-harmony-infer中的 AddKernelInvocation 样例工程位于 ops/ascendc/AddKernelInvocation面向希望在鸿蒙昇腾如 KirinX90 / Kirin9030设备上通过 Ascend C 实现并直调自定义算子的开发者。文章将完整讲解样例的工程结构、Add 算子的 Kernel 实现含多核分块、双缓冲、TPipe/TQue 队列流水、CPU 调测与 NPU 直调两种调用方式以及把样例迁移替换为自有算子时需要修改的全部要点。读完本文你将掌握一条写 Kernel → 编译 → 生成数据 → 运行 → 真值比对的完整算子开发闭环。一、样例工程概览与目录结构AddKernelInvocation是一个算子直调不经过算子注册与推理框架直接在应用程序中调用 Kernel 函数的最小可运行样例固定实现 shape 为8 * 2048的 Add 算子z x y输入输出数据类型为halfFP16。它同时保留了 CPU 侧调测与 NPU 侧运行两条验证路径是理解 Ascend C 算子编程模型的最佳入门工程。其完整目录结构如下原样继承自 README.md├── AddKernelInvocation │ ├── cmake // 编译工程文件含 CCE 编译器适配与 npu 目标构建 │ ├── input // 存放脚本生成的输入数据目录 │ ├── output // 存放算子运行输出数据和真实数据golden的目录 │ ├── scripts │ │ ├── acl.json // acl配置文件本样例为空对象 {} │ │ ├── gen_data.py // 生成输入数据与真值golden数据的脚本 │ │ ├── verify_result.py // 对比算子输出与真值是否一致的验证脚本 │ │── add_custom.cpp // 算子kernel实现暴露给main.cpp调用 │ │── CMakeLists.txt // 编译工程文件 │ │── data_utils.h // 数据读入写出函数ReadFile / WriteFile / CHECK_ACL │ │── main.cpp // 主函数调用算子的应用程序含CPU域及NPU域调用 │ ├── run.sh // 编译运行算子的脚本从源码结构看各文件职责边界清晰add_custom.cpp只负责 Kernel 本身核函数与调用封装main.cpp负责数据准备、设备管理与内核发起data_utils.h提供与设备无关的二进制文件读写scripts/下两个 Python 脚本负责造数据与验结果run.sh将编译、运行、验证串成一条命令。二、Kernel 实现解析CopyIn / Compute / CopyOut 三段式编程2.1 数学表达式与片上执行模型Add 算子的数学表达式为z x y看似简单但它在昇腾设备上的执行模型与 CPU 完全不同Ascend C 提供的矢量计算接口如AscendC::Add操作的元素都是LocalTensor片上 Local Memory 中的张量而输入输出数据本身位于 Global Memory外部存储。因此正确流程是先把输入数据搬运进片上存储再调用计算接口完成加法最后把结果搬出到外部存储。2.2 多核切分与双缓冲参数设计Kernel 在 add_custom.cpp 顶部通过一组constexpr定义了数据规模与并行策略常量值含义TOTAL_LENGTH8 * 2048 16384数据总长度与8*2048的 shape 对应USE_CORE_NUM8使用的 AI Core 数量BLOCK_LENGTHTOTAL_LENGTH / USE_CORE_NUM 2048每个核负责的数据长度TILE_NUM8每个核的数据再切分为 8 个 tileBUFFER_NUM2每个队列的缓冲区个数双缓冲TILE_LENGTHBLOCK_LENGTH / TILE_NUM / BUFFER_NUM 128单次搬入/计算/搬出的数据长度其中TILE_LENGTH 128是由2048 / 8 / 2推导而来由于启用了BUFFER_NUM 2的双缓冲同一份 2048 长度数据被拆成 2 份交替使用每份再切 8 个 tile 流水最终每次处理 128 个元素。整个 Kernel 的循环次数为TILE_NUM * BUFFER_NUM 16次见 add_custom.cpp 的Process()每次循环依次执行 CopyIn → Compute → CopyOut形成搬运与计算重叠的流水线。2.3 KernelAdd 类的四个方法核心类KernelAdd定义于 add_custom.cpp四个方法分别对应1Init绑定全局内存与初始化队列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));通过AscendC::GetBlockIdx()获取当前核的编号将三个全局张量的起始地址各自偏移BLOCK_LENGTH * blockIdx实现多核数据切分——每个核只处理属于自己的 2048 个元素pipe.InitBuffer为输入队列inQueueX、inQueueY与输出队列outQueueZ各分配BUFFER_NUM2块、每块TILE_LENGTH * sizeof(half)字节的片上缓冲区即开启双缓冲。2CopyInGlobal Memory → Local MemoryAscendC::LocalTensorhalf xLocal inQueueX.AllocTensorhalf(); AscendC::LocalTensorhalf yLocal inQueueY.AllocTensorhalf(); AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH); AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH); inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal);从输入队列申请空闲缓冲区用AscendC::DataCopy把 Global Memory 上xGm、yGm的第progress * TILE_LENGTH段搬入再通过EnQue将填好的张量送入队列等待计算。3Compute执行矢量加法AscendC::LocalTensorhalf xLocal inQueueX.DeQuehalf(); AscendC::LocalTensorhalf yLocal inQueueY.DeQuehalf(); AscendC::LocalTensorhalf zLocal outQueueZ.AllocTensorhalf(); AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH); outQueueZ.EnQuehalf(zLocal); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal);从输入队列DeQue取出待计算张量调用矢量计算接口AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH)完成zLocal xLocal yLocal将结果EnQue到输出队列同时FreeTensor归还输入缓冲区以便下一轮 CopyIn 复用——这正是双缓冲流水得以循环运转的关键。4CopyOutLocal Memory → Global MemoryAscendC::LocalTensorhalf zLocal outQueueZ.DeQuehalf(); AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH); outQueueZ.FreeTensor(zLocal);从输出队列取出计算结果用DataCopy写回 Global Memory 上的zGm最后释放输出缓冲区。2.4 核函数入口与对 main 暴露的接口Kernel 入口为extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)见 add_custom.cpp内部依次完成KernelAdd对象的Init与Process。在非 CPU 调测分支#ifndef ASCENDC_CPU_DEBUG下文件还封装了给main.cpp调用的普通 C 函数add_custom_doadd_custom.cppvoid add_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *z) { add_customblockDim, l2ctrl, stream(x, y, z); }这里使用昇腾的内核调用符blockDim, l2ctrl, stream发起 NPU 侧内核执行blockDim指定使用的核数样例为 8stream指定 ACL 流。通过这层包装main.cpp无需感知内核调用符语法只需声明extern void add_custom_do(...)即可完成直调。三、调用实现CPU 调测与 NPU 运行双路径应用程序 main.cpp 通过宏ASCENDC_CPU_DEBUG区分代码逻辑运行于 CPU 侧还是 NPU 侧两条路径在main()内由#ifdef编译期隔离3.1 CPU 侧运行验证ICPU_RUN_KF 调测宏#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); WriteFile(./output/output_z.bin, z, outputByteSize); AscendC::GmFree((void *)x); AscendC::GmFree((void *)y); AscendC::GmFree((void *)z);CPU 侧主要通过 CPU 调测库的ICPU_RUN_KF调测宏在 CPU 上模拟执行核函数先用AscendC::GmAlloc在模拟的 Global Memory 上分配内存ReadFile读入输入二进制SetKernelMode(KernelMode::AIV_MODE)指定矢量核AIV模式再以ICPU_RUN_KF(add_custom, blockDim, x, y, z)触发执行最后WriteFile落盘输出。需要说明的是本样例面向 Kirin 平台main.cpp 中明确注释Kirin暂不支持该编译分支且 run.sh 中亦有Kirin Soc Currently not support cpu!的校验拦截。因此在 KirinX90 / Kirin9030 上实际可用的运行方式是simNPU 仿真。3.2 NPU 侧运行验证ACL 直调链路NPU 分支main.cpp走标准 ACLAscend Computing Language调用链完整流程为aclInit(./scripts/acl.json)初始化 ACL 运行时配置文件本样例为空{}即全部使用默认配置aclrtSetDevice(deviceId)指定设备、aclrtCreateStream(stream)创建流aclrtMallocHost/aclrtMalloc分别申请主机侧与设备侧内存inputByteSize 8 * 2048 * sizeof(uint16_t) 32768字节half即 16 位ReadFile读入input_x.bin、input_y.bin经aclrtMemcpy以ACL_MEMCPY_HOST_TO_DEVICE拷贝到设备调用add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice)发起内核执行随后aclrtSynchronizeStream(stream)同步等待执行完成aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_HOST将结果拷回主机WriteFile写出output_z.bin依次释放设备/主机内存、销毁流、复位设备、aclFinalize()收尾。所有 ACL 调用均通过 data_utils.h 中的CHECK_ACL宏统一检查返回值非ACL_ERROR_NONE时打印出错文件与行号保证错误可定位。四、编译、运行与结果验证4.1 打开样例目录并配置环境变量cd $HOME/ops/AddKernelInvocation说明这里的$HOME需要替换为本仓库根目录即实际路径为ops/ascendc/AddKernelInvocation。export ASCEND_INSTALL_PATH$HOME/Ascend/ascend-toolkit/latest此外从 run.sh 的源码可以看到脚本对环境变量的探测顺序为优先读取ASCEND_TOOLKIT_HOME其次ASCEND_HOME_PATH再次默认路径$HOME/Ascend/ascend-toolkit/latest最后兜底/usr/local/Ascend/ascend-toolkit/latest也支持通过-i/--install-path参数直接传入。4.2 样例执行命令bash run.sh -r [RUN_MODE] -v [SOC_VERSION]两个核心参数说明同时可参考 run.sh 的 getopt 解析逻辑-r / --run-mode运行模式。本样例实际支持cpu / sim / npu三值见 run.sh 的RUN_MODE_LIST校验其中simNPU 仿真是 README 推荐的验证方式cpu模式在 Kirin 平台被显式禁用。-v / --soc-version芯片型号取值为KirinX90或Kirin9030脚本会通过VersionMap校验合法性并据此将CORE_TYPE置为AiCore。-i / --install-path可选直接指定 CANN 安装路径。典型示例bash run.sh -r sim -v KirinX904.3 run.sh 背后编译 → 造数 → 执行 → 验证以sim模式为例脚本实际做了五件事准备仿真环境将LD_LIBRARY_PATH指向$_ASCEND_INSTALL_PATH/x86_64-linux/simulator/${_SOC_VERSION}/lib并设置CAMODEL_LOG_PATH./sim_log见 run.shCMake 编译执行cmake -B build -Dsmoke_testcaseadd -DASCEND_PRODUCT_TYPEKirinX90 -DASCEND_CORE_TYPEAiCore -DASCEND_RUN_MODEsim -DASCEND_INSTALL_PATH...后cmake --build build --target add_sim产出可执行文件add_sim见 run.sh。CMake 侧由 cmake/npu/CMakeLists.txt 将*.cpp以 CCE 语言编译并链接libascendcl生成数据python3 scripts/gen_data.py执行算子运行./add_sim结果验证python3 scripts/verify_result.py output/output_z.bin output/golden.bin。4.4 数据生成与真值比对原理gen_data.py 使用 NumPy 生成 shape 为[8, 2048]、取值范围uniform(1, 100)的 FP16 随机数作为输入input_x、input_y并同步计算真值golden input_x input_y分别落盘到./input/input_x.bin、./input/input_y.bin与./output/golden.bininput_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)verify_result.py 将算子输出output_z.bin与真值golden.bin按 FP16 读入用np.isclose做逐元素近似比较rtol1e-3、atol1e-5容忍equal_nanTrue逐条打印前 100 个不一致元素的下标、期望值、实际值与相对偏差最后以不一致元素占比 ≤ 1e-3作为通过阈值输出test pass。五、替换为自定义算子四步迁移指南如果将该样例工程替换为自己的算子实现需要修改以下内容README 原文建议在run.sh、gen_data.py、main.cpp中搜索注释迁移算子修改点源码中对应的标记分别位于 run.sh、gen_data.py、main.cpp 与 main.cpp1. 替换add_custom.cpp整个文件该文件是核函数实现整个文件都要替换成新算子实现包括新的__global__ __aicore__核函数入口、KernelAdd或其等价类中的Init / CopyIn / Compute / CopyOut逻辑、以及在非 CPU 调测分支下向main.cpp暴露的xxx_do封装函数内部使用内核调用符。若新算子有多个输入、多个输出需同步调整Init中SetGlobalBuffer的参数个数与main侧接口签名。2. 修改run.sh的FILE_NAME在 run.sh 处根据实际算子名修改FILE_NAME属性默认add该变量控制编译获得的二进制名cmake --build build --target ${FILE_NAME}_${RUN_MODE}与后续执行文件./${FILE_NAME}_${RUN_MODE}均依赖它。3. 修改scripts/gen_data.py需自定义生成输入数据的内容包括输入数据的Shape当前为[8, 2048]、数据类型当前为float16需与 Kernel 内half及main.cpp中sizeof(uint16_t)保持一致、算子的输入参数数量当前为 x、y 两个输入并相应修改 golden 的计算公式。4. 修改main.cpp两处迁移算子修改点顶部extern void add_custom_do(...)声明改为调用add_custom.cpp中暴露的新符号见 main.cpp主体读数据与调用逻辑若直调的算子类型改变需结合算子实际输入和gen_data.py同步修改inputByteSize大小、读入文件的数量与顺序、add_custom_do的实参列表见 main.cppCPU 调测分支下同样要调整ReadFile的文件名与ICPU_RUN_KF的入参。完成上述修改后重新执行bash run.sh -r sim -v KirinX90即可完成新算子从编译到验证的全流程。六、小结与延伸阅读AddKernelInvocation样例完整展示了昇腾 Ascend C 算子开发的最小闭环多核切分 TPipe/TQue 队列 双缓冲流水CopyIn → Compute → CopyOut ACL 直调 Python 数据生成与真值比对。理解该样例后读者可以在本仓库 ops/ascendc/src 下找到多个基于相同三段式框架演进的更复杂算子实现如 sobel_custom、rmsNorm_custom、slice_gelu_custom、gather_dequant_int8_custom、quant_matmul_custom 等这些工程在op_hostTiling与op_kernelKernel分离的基础上进一步演示了算子在 ONNX 模型推理链路中的注册与使用方式可作为从直调走向框架集成的下一站参考对应的算子设计文档见 ops/ascendc/docs。若要在鸿蒙端侧完整跑通模型推理可进一步阅读仓库顶层 README.md 与 harmony_infer 目录下的端侧推理示例。【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考