oam-tools 实战AscendC 自定义算子 msprof 采集 AI Core 硬件利用率【免费下载链接】oam-tools本项目为开发者提供故障定位工具包含故障信息收集软硬件信息展示AI core error报错分析等能力提升故障问题定位效率文档可在昇腾社区搜索“故障处理简介”选择社区版。项目地址: https://gitcode.com/cann/oam-tools本文基于仓库 experiment/task-book/msprof_experience_demo/02_api_AscendC/ 中的 AscendC 核函数直调Kernel Launch最小可跑示例展开。文章核心主题是如何手写一个 AscendC 自定义算子用 msprof 采集它在 AI Core 上的 vector/cube 利用率、MTE 搬运占比与 scalar 占比从而判断算子究竟卡在哪个 pipe。读完本文你将掌握一套可复制到任意自研算子的编译 → 采集 → 看数 → 定位瓶颈完整流程。为什么需要给自定义算子做 PMU 采集手写 AscendC kernel 时最关心的问题往往不是能不能跑而是跑得有多快、瓶颈在哪里。一个算子可能同时占用多种硬件单元Vector 单元执行逐元素运算Add、Mul、激活等Cube 单元执行矩阵乘类运算MatMul、卷积MTE搬运引擎负责 GMGlobal Memory与 UBUnified Buffer之间的数据搬入搬出Scalar 单元执行地址计算、循环控制等标量指令。msprof 通过 AI Core PMUPerformance Monitoring Unit硬件计数器把算子在这些单元上的耗时占比统计出来输出到op_summary表。看到vector 只占 4%、搬运占 1/3、标量占主导就能立刻定位算子的 bound 类型而不是靠猜。本 demo 采用昇腾算子开发语言AscendC负载为两个长度 8192 的 fp16 向量逐元素相加element-wise Add并以**核函数直调Kernel Launch**方式运行——不注册算子、不依赖.om离线模型、不依赖任何深度学习框架是 AscendC 的最小可跑单元。相关代码位于 experiment/task-book/msprof_experience_demo/02_api_AscendC/src/add_kernel.cpp。本 demo 属于 oam-tools 仓库中的 msprof 实操任务书task-book一部分与 01_cmdline、03_api_pyAcl、04_pyTorch 一起构成四种采集方式的对照本 demo 的定位是写了算子想看硬件利用率。demo 文件结构与分工文件作用src/add_kernel.cppdevice 侧 AscendC kernelCopyIn → Add → CopyOutsrc/main.cpphost 侧分配显存、拉起 kernel、拷回校验src/CMakeLists.txt用官方ascendc_library编译 device 侧 kernel并链接 host 侧可执行程序build.sh一键编译 →build_run/add_custom_oprun.shmsprof 采集脚本编译产物为build_run/add_custom_op采集输出落在prof_out目录。仓库目录下还有一份 src/add_custom.cpp它是仅示意 AscendC kernel 骨架CopyIn → Compute → CopyOut的未编译样例可对照阅读核心结构。快速跑通编译与采集环境要求Atlas A2 训练系列910B3、CANN 9.1.0 及以上已配置好昇腾开发环境set_env.sh。运行前先source CANN路径/set_env.sh脚本也会自动定位找不到时会给出明确提示。bash build.sh # 编译 bash run.sh 7 # msprof 采集默认 device 7两个脚本均为幂等设计build.sh每次先清空build_run再编译run.sh每次先清空输出目录再重采跑完打印算子聚合结果。build.sh编译链路build.sh 的核心步骤cmake $HERE/src -DASCEND_CANN_PACKAGE_PATH$ASCEND_HOME_PATH \ -DSOC_VERSIONAscend910B3 -DCMAKE_BUILD_TYPERelease cmake --build . -j 4-DASCEND_CANN_PACKAGE_PATHCANN 包路径缺省时回落到$ASCEND_HOME_PATH-DSOC_VERSIONAscend910B3目标芯片型号决定指令集与编译参数。在 src/CMakeLists.txt 中device 侧通过官方ascendc_library编译# 引入官方 AscendC kernel cmake 基础设施提供 ascendc_library set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/ascendc_tools/cmake) if(NOT EXISTS ${ASCENDC_CMAKE_DIR}) set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) endif() include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) # device 侧把 kernel 编成 AscendC 静态库自动生成 aclrtlaunch_add_custom.h ascendc_library(ascendc_kernels STATIC ${CMAKE_CURRENT_SOURCE_DIR}/add_kernel.cpp ) # host 侧可执行 add_executable(add_custom_op ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp) target_link_libraries(add_custom_op PRIVATE ascendc_kernels ascendcl runtime stdc )关键点ascendc_library在编译 device 侧 kernel 的同时会自动生成 host 侧头文件aclrtlaunch_add_custom.h这正是 host 侧 src/main.cpp 里#include aclrtlaunch_add_custom.h的来源——kernel 名add_custom与头文件名一一对应改名时需同步。run.shmsprof 采集命令run.sh 的核心采集命令ASCEND_VISIBLE_DEVICES$DEV msprof --output$OUT \ --ascendclon --runtime-apion --task-timeon --task-memoryon \ --ai-coreon --aic-metricsPipeUtilization --aicpuon --msproftxon \ $BIN参数说明对照仓库 docs/zh/profiling/msprof_cmd/ai_runtime_profile_data.md 中的官方说明--output采集结果输出目录--ascendclon采集 AscendCL API 调用数据--runtime-apion采集 Runtime API 调用数据--task-timeon采集任务下发耗时task 级--task-memoryon采集任务内存占用--ai-coreonAI Core 数据采集开关采集硬件 PMU 指标的前提--task-time为 on 时默认开启--aic-metricsPipeUtilization采集计算类和搬运类指令耗时和占比即本 demo 要看的aiv_vec_ratio、aiv_mte2_ratio等 pipe 利用率指标。它是所有昇腾训练/推理系列产品含 Atlas A2 系列、Ascend 950 系列的通用推荐取值PipeUtilization也可统计*_icache_miss_rate--aicpuon采集 AI CPUAICPU任务数据--msproftxon开启 msproftx 打点。--aic-metrics还支持更细粒度的指标如ArithmeticUtilization算数单元利用率、Memory、L2Cache、MemoryBandwidth等以及Custom:0x49,0x8,...自定义寄存器最多 8 个。本 demo 只关心 pipe 占比用PipeUtilization即可。采集完成后脚本自动打印算子聚合结果cat $OUT/PROF_*/mindstudio_profiler_output/op_statistic_*.csv看懂代码从 host 到 device 的完整链路device 侧CopyIn → Add → CopyOutsrc/add_kernel.cpp 定义了一个典型的 AscendC 数据流constexpr int32_t TOTAL_LENGTH 8192; // 总元素数 constexpr int32_t USE_CORE_NUM 8; // 用 8 个核 constexpr int32_t BLOCK_LENGTH TOTAL_LENGTH / USE_CORE_NUM; // 每核处理量 constexpr int32_t TILE_NUM 8; // 每核切 8 片流水 constexpr int32_t BUFFER_NUM 2; // double buffer constexpr int32_t TILE_LENGTH BLOCK_LENGTH / TILE_NUM / BUFFER_NUM;在Init()中三个 GlobalTensor 按GetBlockIdx()切分到各核TPipe与三个队列TQueQuePosition::VECIN, BUFFER_NUM、TQueQuePosition::VECOUT, BUFFER_NUM构成double buffer乒乓缓冲——搬入与计算在同一时刻重叠进行这是让 MTE 与 vector 单元并行的关键。主循环为Process()__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); } }三段式流水分别对应CopyInDataCopy把 GM 数据搬入 UBinQueueX/inQueueY入队EnQueCompute出队DeQue拿到输入调用 vector 指令Add(zLocal, xLocal, yLocal, TILE_LENGTH)逐元素相加结果写入输出队列CopyOutDataCopy把 UB 结果搬回 GMzGm。最后通过extern C __global__ __aicore__导出入口extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelAdd op; op.Init(x, y, z); op.Process(); }host 侧分配显存 → 拉起 kernel → 拷回校验src/main.cpp 用纯 AscendCL API 完成 host 侧逻辑aclInit(nullptr)→aclrtSetDevice(0)→aclrtCreateStream初始化运行时并建流aclrtMallocHostaclrtMalloc分配 host 与 device 内存ACL_MEM_MALLOC_HUGE_FIRST优先申请大页内存初始化为 fp16 的 1.00x3C00aclrtMemcpy两次 H2D 拷入 x、y用ACLRT_LAUNCH_KERNEL宏直调 kernel循环 20 次拉起并aclrtSynchronizeStream同步——多次执行为 msprof 提供稳定的 PMU 采样样本避免单次执行样本过少导致统计偏差for (int it 0; it 20; it) { ACLRT_LAUNCH_KERNEL(add_custom)(BLOCK_DIM, stream, xD, yD, zD); } CHECK(aclrtSynchronizeStream(stream));D2H 拷回结果并校验z[0]与z[8191]应为 0x4000即 2.0验证 kernel 正确性依次aclrtFree、aclrtDestroyStream、aclrtResetDevice、aclFinalize释放资源。所有 ACL 调用都通过CHECK宏断言ACL_SUCCESS任一失败即打印[FAIL]并退出——采集环境出问题时能第一时间暴露而不是产出脏数据。预期结果910B3 示例与核心洞察在 910B3 上实跑后op_summary单算子 PMU 画像中逐元素 Add 的典型指标如下PMU 指标值解读aiv_vec_ratio~0.04vector 计算只占 4%——数据太少算得太快aiv_scalar_ratio~0.65标量指令占主导地址计算/循环控制aiv_mte2_ratio~0.33从 GM 搬入占 1/3cube_utilization(%)0element-wise 不用 cube符合预期核心洞察Add 这类 element-wise 算子是scalar/搬运 bound标量与搬运受限真正 vector 计算只占 4%。与之形成鲜明对比的是 MatMul 这类矩阵乘算子——cube boundcube_utilization可达 60% 以上。这两种完全不同的瓶颈画像正是 msprof 采集 AscendC 算子最大的价值用硬件计数器替代拍脑袋估算直接回答算子到底卡在哪个 pipe。在解读op_summary时还需注意aiv_vec_ratio低不代表算子写得差而是负载特性使然——数据量小、vector 单元空转等待搬运与标量控制若发现aiv_mte2_ratio很高说明搬运成为瓶颈可考虑增大 tiling、加深流水更多 buffer 数以提升搬运与计算重叠度若aiv_scalar_ratio异常偏高可审视循环内是否有过多地址计算或分支判断把--aic-metricsPipeUtilization换成L2Cache、MemoryBandwidth等指标可进一步定位 L2 命中率与带宽瓶颈仓库 docs/zh/profiling/msprof_cmd/ai_runtime_profile_data.md 有各指标的完整取值范围与适用产品说明。如何用到你自己的算子把本 demo 的 Add 换成任意自研算子只需改三处src/add_kernel.cpp改写Compute()中的计算逻辑必要时同步调整Init()里的 buffer 配置、队列位置与类型例如矩阵乘需要QuePosition::CUBEIN/CUBEOUTsrc/main.cpp按新 kernel 的签名修改ACLRT_LAUNCH_KERNEL的 launch 名与输入输出参数src/CMakeLists.txt把ascendc_library的源文件列表换成你的 kernel 文件。然后重新执行bash build.sh bash run.sh 7若你的算子编译/采集环境不同注意同步确认SOC_VERSION在build.sh中为 Ascend910B3与 CANN 版本是否匹配。小结本 demo 给出了一个端到端的最小闭环AscendC kernelCopyIn→Compute→CopyOut→ascendc_library编译 → host 侧 Kernel Launch 直调 → msprof--ai-coreon --aic-metricsPipeUtilization采集 →op_summary解读。对 Add 这类 element-wise 算子实测画像呈现典型的 scalar/搬运 boundaiv_scalar_ratio约 0.65、aiv_mte2_ratio约 0.33、aiv_vec_ratio仅约 0.04、cube_utilization为 0把这套流程套用到你自己的算子上即可量化每种硬件单元的占用快速判断算子瓶颈类型。更完整的 msprof 参数与指标说明可继续阅读仓库 docs/zh/profiling/msprof_cmd/msprof_cmd.md 与 docs/zh/profiling/README.md。【免费下载链接】oam-tools本项目为开发者提供故障定位工具包含故障信息收集软硬件信息展示AI core error报错分析等能力提升故障问题定位效率文档可在昇腾社区搜索“故障处理简介”选择社区版。项目地址: https://gitcode.com/cann/oam-tools创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考