msprof 性能采集实战指南：在 CANN oam-tools 中用四种方式采集昇腾 NPU 算子性能数据

msprof 性能采集实战指南：在 CANN oam-tools 中用四种方式采集昇腾 NPU 算子性能数据

【免费下载链接】oam-tools本项目为开发者提供故障定位工具，包含故障信息收集，软硬件信息展示，AI core error报错分析等能力，提升故障问题定位效率，文档可在昇腾社区搜索“故障处理简介”（选择社区版）。项目地址: https://gitcode.com/cann/oam-tools

导读

本文基于experiment/task-book/msprof_experience_demo示例仓库，系统讲解在昇腾 NPU 上使用 msprof 进行性能数据采集的四种实战方式：msprof 命令行黑盒采集、AscendC 自定义算子采集、pyACL 离线模型推理采集、torch_npu.profilerAPI 白盒插桩采集。全文以最简 TinyMLP 为统一负载，逐一演示每种方式的原理、代码结构、运行命令与输出解读，并结合仓库源码剖析 msprof 关键采集参数（--ai-core、--aic-metrics=PipeUtilization等）的底层作用，帮助读者掌握"跑通 msprof、看懂性能数据、定位算子热点"的完整链路，并能将示例负载替换为自己的模型或算子。

一、Demo 概览：四种采集方式的定位

msprof_experience_demo是一套 msprof 上手示例，用同一个最简 TinyMLP 负载（4 层 Linear + GELU，输入[32, 1024]）演示昇腾 NPU 上四种性能数据采集方式。其设计意图非常明确：把 msprof 跑通、看懂输出，真正使用时把负载换成你自己的模型即可（源码中统一用← 换成你的模型注释标注替换点）。

四种方式及适用场景如下表（来自顶层 README.md）：

四种脚本均可实跑；运行后性能数据落在各方式run.sh指定的输出目录，仓库本身不附带实采数据。

目录结构

experiment/task-book/msprof_experience_demo/ ├── README.md ├── 01_cmdline/ src/model.py + run.sh 命令行黑盒采集 ├── 02_api_AscendC/ src/(add_kernel.cpp + main.cpp + CMakeLists.txt) + build.sh + run.sh 算子编译+采集 ├── 03_api_pyAcl/ src/(export_onnx.py + infer.py) + build_model.sh + run.sh ONNX→om→推理 └── 04_pyTorch/ src/model_with_profiler.py + run.sh profiler API 插桩

每个子目录都有独立 README，讲清该方式的细节以及"如何用到你的模型"。

二、环境准备

示例运行环境如下（仓库实测环境）：

• Atlas A2 训练系列（910B3），CANN 9.1.0，torch_npu 2.7.1，Python 3.12
• 跑各方式前先执行source <CANN路径>/set_env.sh（脚本会自动定位 CANN 环境，找不到会提示报错）

启动命令统一为：

cd 01_cmdline && bash run.sh 7 # 命令行 cd 02_api_AscendC && bash build.sh && bash run.sh 7 # AscendC(先编译) cd 03_api_pyAcl && bash build_model.sh && bash run.sh 7 # pyACL(先转 om) cd 04_pyTorch && bash run.sh 7 # PyTorch API

其中参数7是 device 号，可按实际环境换成其他卡号。脚本均为幂等设计：每次运行先清空输出目录再重新采集，跑完自动打印 Top 算子。

预期输出文件

采集结果落在各方式run.sh指定的输出目录，核心是可读的 CSV/JSON 文件：

示例 TinyMLP 的实测结论：MatMulV2 占约 78%（主热点）；step_trace 显示每步"等待"是"计算"的约 60 倍，说明模型太小、瓶颈在 host 下发。

三、公共负载：最简 TinyMLP

01_cmdline/src/model.py（源码）与04_pyTorch/src/model_with_profiler.py（源码）使用同一个 TinyMLP：3 组Linear(1024,1024) + GELU再接一层Linear(1024,1024)，输入[32, 1024]。

def build_model(): # ← 换成你的模型 return torch.nn.Sequential( torch.nn.Linear(1024, 1024), torch.nn.GELU(), torch.nn.Linear(1024, 1024), torch.nn.GELU(), torch.nn.Linear(1024, 1024), torch.nn.GELU(), torch.nn.Linear(1024, 1024), )

设计上刻意保持"小、纯、单一算子"：固定 shape、warmup 3 次 + 测量 20 次，目的是给 msprof 一个稳定、可复现、MatMul 主导的负载，便于对比四种采集方式的输出是否自洽。

四、方式一：msprof 命令行（CLI）黑盒采集

原理

msprof 作为父进程拉起你的应用，全程旁路采集，不改一行模型代码。适用于"拿到一个能跑的黑盒应用，想快速看一眼性能"的场景。注意一个关键细节：卡号必须写在 msprof 命令之前，即ASCEND_VISIBLE_DEVICES=$DEV msprof ...，因为 msprof 是父进程，需要在其初始化前让环境变量生效，否则会落到 device 0。

采集脚本与参数逐项解读

run.sh 的核心采集命令：

ASCEND_VISIBLE_DEVICES=$DEV msprof --output="$OUT" \ --ascendcl=on --runtime-api=on --task-time=on --task-memory=on \ --ai-core=on --aic-metrics=PipeUtilization --aicpu=on --msproftx=on \ python3 "$HERE/src/model.py"

各参数作用如下（可与官方参数文档 ai_runtime_profile_data.md 对照阅读）：

跑完脚本自动cat输出op_statistic_*.csv打印算子耗时 Top。

预期结果（910B3 示例）

算子耗时 Top（op_statistic_*.csv）：

MatMul 占主导，符合全连接网络预期。op_summary_*.csv含完整 AI Core PMU 指标，可进一步定位单算子瓶颈。

如何用到你的模型

零侵入：把run.sh最后一行的python3 "$HERE/src/model.py"换成你自己的启动命令（任何在 NPU 上有真实计算的程序都行），其余 msprof 参数不动即可。

五、方式二：AscendC 自定义算子 + msprof 采集

这种方式采什么

手写一个 AscendC kernel（昇腾的算子开发语言），msprof 采集它在 AI Core 上的vector/cube 利用率、MTE 搬运占比、scalar 占比，判断算子卡在哪个 pipe。Demo 使用最小可跑单元——核函数直调（Kernel Launch），负载是两个长度 8192 的 fp16 向量逐元素相加。

文件构成

源码剖析

device 侧 kernel（add_kernel.cpp）：经典的三段式流水结构。

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()：通过xGm.SetGlobalBuffer(...)按GetBlockIdx()切分全局内存，每个核处理自己的一段；pipe.InitBuffer为输入/输出队列各申请BUFFER_NUM片 buffer。
• Process()：循环TILE_NUM * BUFFER_NUM次，交替执行CopyIn(i)（GM → Local，即 MTE 搬运）、Compute(i)（vector 单元执行Add）、CopyOut(i)（Local → GM）。双 buffer 机制让搬运与计算重叠，是衡量流水线利用率的关键结构。
• 入口extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)不做算子注册、不依赖 .om 或框架，是 AscendC 最小可跑单元。

host 侧启动程序（main.cpp）：aclInit→aclrtSetDevice→aclrtCreateStream→ 分配 host/device 内存 →aclrtMemcpy拷贝输入 → 通过ACLRT_LAUNCH_KERNEL(add_custom)(BLOCK_DIM, stream, xD, yD, zD)循环拉起 kernel 20 次（warmup + 多次拉起给 msprof 稳定样本）→ 同步后拷回校验（fp16 的 1+1=2，期望输出 0x4000）→ 逆序释放资源。

CMake 配置（CMakeLists.txt）：引入 CANN 包内tools/ascendc_tools/cmake（或compiler/tikcpp/ascendc_kernel_cmake）下的官方ascendc.cmake，用ascendc_library把add_kernel.cpp编成静态库并自动生成aclrtlaunch_add_custom.h（main.cpp 即 include 此头文件），host 侧链接ascendcl、runtime。

编译与采集

bash build.sh # 编译，SOC=Ascend910B3，产物 build_run/add_custom_op bash run.sh 7 # msprof 采集（默认 device 7）

注意run.sh中先检查可执行文件是否存在，并export LD_LIBRARY_PATH="$ASCEND_HOME_PATH/aarch64-linux/lib64:$LD_LIBRARY_PATH"，msprof 参数与方式一完全一致（--ai-core=on --aic-metrics=PipeUtilization ...）。

预期结果（910B3 示例）

单算子 PMU（op_summary，逐元素 Add 的典型画像）：

核心洞察：Add 这种 element-wise 算子是scalar/搬运 bound，真正 vector 计算只占 4%；与 MatMul（cube bound，cube_utilization60%+）形成鲜明对比——这正是 msprof 采集 AscendC 算子的价值：用 PMU 量化算子到底卡在哪个 pipe，指导 kernel 优化（如减少标量开销、增大数据块、优化搬运）。

如何用到你的算子

改src/add_kernel.cpp的Compute()计算逻辑、src/main.cpp的输入输出与 kernel 名、src/CMakeLists.txt的源文件，然后bash build.sh && bash run.sh 7。

六、方式三：pyACL 加载 .om 离线模型推理 + msprof 采集

这种方式采什么

pyACL 是 CANN 的 Python 推理接口，面向离线模型（.om）部署场景：模型先用 ATC 从 ONNX 转成 .om，再用acl.mdl.execute推理。Demo 用与 01/04 相同的 TinyMLP，走 ONNX → .om 转换链后用 pyACL 推理，便于横向对比部署链路与框架推理的差异。

文件构成与完整链路

完整链路：export_onnx.py（TinyMLP→ONNX）→atc（ONNX→.om）→infer.py（pyACL 推理）。

ATC 转换

build_model.sh 中的核心命令：

python3 "$HERE/src/export_onnx.py" # TinyMLP → tiny_mlp.onnx atc --model=tiny_mlp.onnx --framework=5 --output=tiny_mlp \ --soc_version=Ascend910B3 --input_shape="x:32,1024"

其中--framework=5表示 ONNX 框架，--soc_version=Ascend910B3指定芯片型号，--input_shape="x:32,1024"指定输入名与 shape。

pyACL 推理全链路

infer.py 遵循 pyACL 标准调用顺序：

1. acl.init()→acl.rt.set_device(DEVICE_ID)→acl.rt.create_context(DEVICE_ID)
2. acl.mdl.load_from_file(om_path)加载 .om 模型，acl.mdl.create_desc()+acl.mdl.get_desc()获取模型描述
3. 准备输入：host 侧np.ones((32,1024), dtype=np.float32)→acl.rt.malloc分配 device 内存 →acl.rt.memcpyH2D →acl.mdl.create_dataset()+acl.create_data_buffer+acl.mdl.add_dataset_buffer组装 input dataset
4. 按acl.mdl.get_output_size_by_index(model_desc, 0)查询输出大小并分配输出 dataset
5. 执行推理 ×20（warmup 3 + 测量 17），acl.mdl.execute(model_id, in_ds, out_ds)——msprof 在此采集算子
6. 拷回输出校验：acl.rt.malloc_host→acl.rt.memcpyD2H →acl.util.ptr_to_bytes转 numpy
7. 逆序清理资源（dataset → buffer → device 内存 → desc → unload → context → reset_device → finalize）

预期结果（910B3 示例）与关键洞察

算子聚合（op_statistic）——与 01/04 对比出现关键差异：

核心洞察：pyACL 路径凭空多出Cast 算子（占 28.6%），而 PyTorch 路径没有。原因是 ATC 转 .om 默认按 fp16 优化，而模型 IO 是 fp32，于是自动插入 fp32↔fp16 转换。这是离线部署相对框架推理的典型代价——可在 ATC 阶段让 IO 也走 fp16 以消除这些 Cast。

如何用到你的模型

替换src/export_onnx.py的模型与输入（或直接拿已有 ONNX 改build_model.sh的 atc 命令），改src/infer.py的输入 shape，然后bash build_model.sh && bash run.sh 7。

七、方式四：torch_npu.profiler API 白盒采集

这种方式是什么

在代码里用torch_npu.profiler.profile插桩，白盒采集。相比 CLI（01），API 能精确圈定第 N~M step、拿到 aten op 级耗时与 Python 调用栈，并额外产出 CLI 没有的step_trace_time.csv（每 step Computing/Free 占比）。

采集配置详解

model_with_profiler.py 中把可调"旋钮"单独列出：

level = ProfilerLevel.Level1 # 采集层级，Level1 含 AI Core PMU；想更轻量可降 Level0 pmu = AiCMetrics.PipeUtilization # AI Core 指标口径，看各计算流水线占用 acts = [ProfilerActivity.CPU, ProfilerActivity.NPU] # CPU+NPU 都要，才能看清"下发 vs 执行"的 gap exp_config_cls = getattr(torch_npu.profiler, "_ExperimentalConfig") trace_cfg = exp_config_cls( profiler_level=level, aic_metrics=pmu, data_simplification=False, # 生产场景可设 True 省空间 ) window = schedule(wait=0, warmup=WARMUP_STEPS, active=ACTIVE_STEPS, repeat=1) profiler = profile( activities=acts, schedule=window, on_trace_ready=tensorboard_trace_handler(out_dir), experimental_config=trace_cfg, record_shapes=True, # 记录算子 input shape with_stack=True, # 记录 Python 调用栈 )

主循环中每次迭代末尾必须调用prof.step()通知 profiler 进入下一 step，采集区间由schedule(wait/warmup/active)圈定（Demo 为 warmup 3 + active 5）。源码通过getattr按名取用_ExperimentalConfig，避免直接书写受保护成员。

实测结果（910B3 示例）

算子聚合（op_statistic.csv）——与 CLI 高度一致，验证两种方式自洽：

step 拆分（step_trace_time.csv，API 独有）：

关键洞察：NPU 实际计算每 step 仅约 35us，但 Free（等待 host 下发）高达约 2000us，是典型的host bound——TinyMLP 太小，host 下发开销完全盖过了计算。这正是 API 模式相对 CLI 的价值：CLI 只告诉你"MatMul 占 78%"，API 进一步告诉你"整个 step 99% 时间在等 host"。

host 下发 Top（api_statistic.csv）

如何用到你的模型

1. 替换src/model_with_profiler.py里的build_model()和输入x（均标了← 换成你的模型注释）。
2. 训练循环则把forward + backward + optimizer.step()放进with prof:块内，每 step 末尾调prof.step()，用schedule(wait/warmup/active)圈定要采的 step。
3. 采集配置（make_profiler那段）和run.sh不用改，直接bash run.sh 7。

换成自己的模型后，重点看step_trace_time.csv的 Computing:Free——它告诉你是算力不够（Computing 高）还是 host 下发拖后腿（Free 高）。

八、四种方式对比与选型建议

实践建议：

• 快速摸底一个能跑的黑盒程序→ 01 CLI，改一行命令即可；
• 写了自己的 AscendC 算子、想看硬件利用率→ 02，用 PMU 判断卡在哪个 pipe；
• 离线部署 .om 模型→ 03，同时留意 ATC 精度优化引入的 Cast 开销；
• PyTorch 场景、想定位是算力不够还是 host 下发拖后腿→ 04，step_trace_time.csv直接给出 Computing:Free 比例。

九、输出文件深度解读

• op_statistic.csv：算子类型级聚合（Count、Total、Ratio），用于回答"谁是热点算子"；四种方式均可对比该文件验证采集自洽性（01 与 04 的 MatMulV2 占比分别为 77.88% 与 78.30%）。
• op_summary_*.csv：单算子级明细 + AI Core PMU，如方式二中aiv_vec_ratio/aiv_scalar_ratio/aiv_mte2_ratio/cube_utilization(%)用于定位计算单元瓶颈。
• step_trace_time.csv：API 独有，每 step 的 Computing/Free 拆分，用于判断 host bound 还是 compute bound。
• trace_view.json：完整时间线，可拖进chrome://tracing（或 Perfetto UI）逐算子查看执行窗口。

十、延伸阅读

• 四种方式各自的 01_cmdline/README.md、02_api_AscendC/README.md、03_api_pyAcl/README.md、04_pyTorch/README.md
• 官方 msprof 命令行采集参数说明：ai_runtime_profile_data.md、msprof_cmd.md
• 采集选项与进阶配置：profiling_options_parameter.md、PROFILING_OPTIONS.md、PROFILING_MODE.md
• msprof 采集器源码实现：collector/dvvp/msprof、collector/dvvp/profimpl，可深入研读 AI Core 数据采集与 PMU 指标计算的底层逻辑

【免费下载链接】oam-tools本项目为开发者提供故障定位工具，包含故障信息收集，软硬件信息展示，AI core error报错分析等能力，提升故障问题定位效率，文档可在昇腾社区搜索“故障处理简介”（选择社区版）。项目地址: https://gitcode.com/cann/oam-tools