目录结构介绍【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills├── kirin_add_template │ ├── cmake // 编译工程文件 │ │ ├── Modules/ // CCE编译器检测模块通用无需修改 │ │ └── npu/CMakeLists.txt // NPU编译配置通用无需修改 │ ├── input // 存放脚本生成的输入数据目录 │ ├── output // 存放算子运行输出数据和真值数据的目录 │ ├── scripts │ │ ├── acl.json // ACL初始化配置默认空JSON │ │ ├── gen_data.py // 生成输入数据和golden真值需修改 │ │ └── verify_result.py // 输出与golden对比验证需修改 │ ├── add_custom.cpp // 算子核函数实现需替换 │ ├── data_utils.h // 通用文件读写/日志工具无需修改 │ ├── main.cpp // Host侧主程序需修改 │ ├── CMakeLists.txt // 顶层CMake配置无需修改 │ └── run.sh // 一键编译执行验证脚本需修改代码实现介绍本调用样例中实现的是固定shape为8*2048的Add算子。kernel实现Add算子的数学表达式为z x y计算逻辑是Ascend C提供的矢量计算接口的操作元素都为LocalTensor输入数据需要先搬运进片上存储然后使用计算接口完成两个输入参数相加得到最终结果再搬出到外部存储上。Add算子的实现流程分为3个基本任务CopyInComputeCopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm和yGm搬运到Local Memory分别存储在xLocal、yLocalCompute任务负责对xLocal、yLocal执行加法操作计算结果存储在zLocal中CopyOut任务负责将输出数据从zLocal搬运至Global Memory上的输出Tensor zGm中。具体请参考add_custom.cpp。调用实现CPU侧运行验证主要通过ICPU_RUN_KF CPU调测宏等CPU调测库提供的接口来完成NPU侧运行验证主要通过使用内核调用符来完成。应用程序通过ASCENDC_CPU_DEBUG 宏区分代码逻辑运行于CPU侧还是NPU侧。运行样例算子打开样例目录cd $HOME/ops/kirin_add_template配置环境变量这里的$HOME需要替换为本仓根目录export ASCEND_INSTALL_PATH$HOME/Ascend/cann-mobile/cann-8.5.0样例执行bash run.sh -r [RUN_MODE] -v [SOC_VERSION]RUN_MODE编译方式simNPU仿真。SOC_VERSIONKirinX90 或者 Kirin9030。示例如下。bash run.sh -r sim -v Kirin9030run.sh的完整执行流程为解析参数通过-vSoC版本、-r运行模式sim/cpu/npu、-iCANN安装路径解析命令行参数确定CANN路径按优先级从ASCEND_TOOLKIT_HOME→ASCEND_HOME_PATH→ 用户指定路径 → 默认路径查找校验参数验证SoC版本必须是 KirinX90/Kirin9030运行模式必须是 sim/cpu/npuKirin SoC不支持cpu模式配置仿真环境如果是sim模式设置模拟器库路径LD_LIBRARY_PATH和日志目录编译清理build目录通过CMake配置并编译生成可执行文件生成测试数据用gen_data.py生成输入和golden数据执行运行编译出的可执行文件在模拟器上执行算子验证结果用verify_result.py对比输出与golden数据确认误差在容忍范围内清理删除仿真日志和vcd文件基于本模板开发新自定义算子以下以开发一个mul乘法算子为例说明如何将本模板改造为新算子工程。步骤1复制模板并重命名cp -r kirin_add_template mul_template cd mul_template rm -rf build add_sim *_cpu *_npu cceprint npuchk *.vcd sim_log input/*.bin output/*.bin步骤2修改run.sh— FILE_NAMErun.sh:27中的FILE_NAME控制编译产物名称和CMake target名。改为新算子名FILE_NAMEmul # 原: FILE_NAMEadd步骤3替换add_custom.cpp— 核函数实现这是最核心的修改整个文件替换为新算子实现。需要修改的内容计算常量行19-24根据新算子的数据总量、核心数、分块策略调整constexpr int32_t TOTAL_LENGTH 8 * 2048; constexpr int32_t USE_CORE_NUM 1; constexpr int32_t BLOCK_LENGTH TOTAL_LENGTH / USE_CORE_NUM; constexpr int32_t TILE_NUM 8; constexpr int32_t BUFFER_NUM 2; constexpr int32_t TILE_LENGTH BLOCK_LENGTH / TILE_NUM / BUFFER_NUM;Kernel类名和逻辑替换类名修改Init/Process/CopyIn/Compute/CopyOutInit调整 GlobalTensor 数量和类型Compute将AscendC::Add替换为目标计算接口例如AscendC::MulCopyIn/CopyOut调整输入/输出队列数量和搬运逻辑核函数入口行84-88改名为新算子名调整参数列表extern C __global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelMul op; op.Init(x, y, z); op.Process(); }Host调用桥接函数行94-97改函数名与核函数名一致void mul_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *z) { mul_customblockDim, l2ctrl, stream(x, y, z); }步骤4修改main.cpp— Host侧调用接口extern声明行21/24改为新算子名extern void mul_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z); // CPU分支: extern C __global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z);数据大小行30-31根据新算子的输入/输出shape和数据类型调整size_t inputByteSize 8 * 2048 * sizeof(uint16_t); // half uint16_t size_t outputByteSize 8 * 2048 * sizeof(uint16_t);调用点行75改函数名参数与核函数保持一致mul_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice);输入数据读取行68-69如果输入参数数量/名称变化需同步调整ReadFile的文件名步骤5修改scripts/gen_data.py— 生成输入和真值根据新算子的数学定义修改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) # 改 Add 为 Mul input_x.tofile(./input/input_x.bin) input_y.tofile(./input/input_y.bin) golden.tofile(./output/golden.bin)需要修改的内容Shape如果算子shape不同修改[8, 2048]及数据类型np.float16/np.float32/np.int32等输入参数数量增减输入文件如3输入算子需要额外生成input_w.bin计算公式将改为目标运算输出文件名如有多输出需生成多个 golden 文件步骤6修改scripts/verify_result.py— 验证逻辑根据输出数据类型调整# 如输出为 float32: output np.fromfile(output, dtypenp.float32).reshape(-1) golden np.fromfile(golden, dtypenp.float32).reshape(-1) # 调整容差float32精度更高时可用更小容差 relative_tol 1e-5 absolute_tol 1e-8 error_tol 1e-5无需修改的文件文件原因data_utils.h通用工具函数与算子无关CMakeLists.txt顶层只做add_subdirectory(cmake/npu)无算子名cmake/npu/CMakeLists.txt使用${smoke_testcase}变量由CMake参数传入cmake/Modules/*CCE编译器检测模块通用scripts/acl.jsonACL配置默认为空{}常见扩展场景的修改要点场景需修改的文件改数据类型half→floatadd_custom.cppGlobalTensor/LocalTensor类型、main.cppsizeof、gen_data.pydtype、verify_result.pydtype容差改Shapeadd_custom.cppTOTAL_LENGTH等常量、main.cppinputByteSize、gen_data.pyshape多输入3个以上add_custom.cpp增加队列和GlobalTensor、main.cpp增加malloc/read/memcpy、gen_data.py增加输入文件多输出add_custom.cpp增加输出队列、main.cpp增加mallocWriteFile、gen_data.py增加golden文件、verify_result.py多输出对比多核并行add_custom.cpp调整USE_CORE_NUM和GetBlockIdx逻辑、main.cpp调整blockDim关键概念速查核函数extern C __global__ __aicore__修饰的函数是NPU上的执行入口_do桥接函数在核函数实现文件中用调用符包装核函数供main.cpp的Host侧调用ASCENDC_CPU_DEBUG宏区分CPU调测分支和NPU/仿真分支Kirin SoC目前仅支持NPU/仿真分支双缓冲BUFFER_NUM2通过队列深度为2实现流水线一个buffer搬运时另一个buffer计算【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考