Ascend Operator Performance Tuning

对可运行且精度通过的昇腾算子执行“环境与版本对齐 → 编译 → 基线 → msOpProf 采集 → 瓶颈判定 → 单变量优化 → 前后对比”。所有性能结论来自目标环境实测；知识库案例只提供候选机制，不替代当前算子的 profiling 证据。

使用本技能时，咨询和执行任务都必须先完成资料查询：在代码仓 https://gitcode.com/Ascend/msopprof 查询 msopprof 与 msopprof simulator 使用资料。下文及 references 中的参数、能力边界和命令模板按返回的相关正文核验后使用；查询失败时查找本地相关工具或本地资料参考用法回答和命令执行。

任务边界

• 已有源码、构建入口和测试入口，用户要求性能分析或调优：执行本流程。
• 精度未通过、运行崩溃或 kernel 挂起：停止性能调优，转交精度或运行时调试流程。
• 从零开发、CUDA 迁移或接口设计：先完成开发流程，算子可运行后再使用本 Skill。
• 只要求解释 profiling 指标：只执行相应诊断，不修改源码。

必须解析的输入

从用户信息、工程和当前环境依次获取；只有缺失字段会改变执行路径时才询问。

字段	获取位置	缺失时处理
工程根目录、算子名、DSL	用户输入、构建文件、源码后缀	无法识别构建入口或 DSL 时询问
kernel 清单	注册信息、运行日志、OpBasicInfo.csv	多 kernel 场景必须形成完整清单
CANN/driver/firmware、SoC	版本文件、npu-smi、ACL/torch_npu	不明确时只做离线审查，不声称上板结果
shape/dtype/format/属性	测试配置、调用入口	禁止用临时 smoke shape 替代目标用例
正确性命令与容差	工程现有测试	禁止自创业务容差
性能命令、warmup、repeat	工程 benchmark	无等价标杆时只做前后自对比

渐进式披露规则

不要一次性加载整个 references/：

1. 识别 DSL 和算子类型后，只读取匹配的 compile-*.md。
2. 准备采集时读取 msOpProf 采集指南。
3. 根据已采集字段选择一个或多个 diagnose-*.md；没有证据时不得先加载案例猜瓶颈。
4. Bound 确认后读取 优化技术索引，再加载命中的单个技术文件。
5. 最后读取 案例路由，按 DSL、算子类型、SoC 和 Bound 选择最多三个同型案例。

执行流程

Step 0：建立隔离产物目录

不覆盖用户基线。源码修改放在用户允许的分支或副本；运行产物写入被忽略的目录：

tuning_artifacts/<operator>/
├── environment.txt
├── scope.json
├── baseline/
├── round-01/
└── comparison.md

scope.json 至少记录：operator、dsl、repo_commit、cann_version、soc、device_id、shapes、dtypes、formats、kernel_names、三类执行命令、warmup、repeat 和 baseline_kind。

Step 1：环境与版本对齐

bash <skill_dir>/scripts/env_check.sh --mode <board|sim> --cann-path <cann-path>

把输出保存到 environment.txt，并确认：

1. 源码 branch/tag 与 CANN/OPP 版本兼容；不匹配时切换对应 release 或记录 source_mismatch=true。
2. A2/A3/A5 编译目标来自目标仓文档、本机工具链或 build.sh --help，不按机器昵称猜测。
3. 使用 msprof op --help 或 msopprof --help 记录当前版本真实支持的参数、指标和 simulator/MC2 能力。
4. 选择健康且空闲的 NPU，记录设备 ID、频率和同卡负载。
5. 通信/多 rank 算子：rankNum ≤ 可用空闲设备数 在本步判定，不满足直接标记 BLOCKED（host 侧 deviceId = rankId 的样例每 rank 独占一个物理设备）。
6. Triton-Ascend 算子：可用性门禁是"最小 triton kernel 上板跑通"（vector add 级），不是 import triton 成功；版本配套与 207000 排查见 compile-triton.md §2.7。
7. 数据/校验脚本的 Python 依赖：逐个 python3 -c "import <mod>" 核对 gen_data/verify 脚本的真实 import（常见缺口：torch、en_dtypes、ml_dtypes、pandas），缺失时在产物目录建隔离 venv，不污染系统环境。

缺 simulator 但执行上板路径时为非阻塞。没有可用 NPU 时只完成编译、路由和命令生成，不生成性能数字。

Step 2：按 DSL 编译并通过精度门禁

DSL/框架	唯一入口
Ascend C / CANN 官方算子仓	compile/compile-ascendc.md
CATLASS	compile/compile-catlass.md
Triton-Ascend	compile/compile-triton.md
TileLang-Ascend	compile/compile-tilelang.md
PyPTO	compile/compile-pypto.md
SHMEM / MC2	compile/compile-shmem.md

保存完整构建命令、退出码和产物位置。profiling 构建保持 Release 优化；需要源码映射时只增加目标编译器支持的调试信息。运行工程原有正确性全集并记录容差、最大绝对/相对误差和失败用例。任一目标用例失败即停止性能调优。

精度保障机制（全流程三道门禁 + 校验链审计）

精度不是"最后复测一次"，而是三次强制门禁，任何一道不过即停止或回滚：

1. 进入门禁（Step 2）：基线建立前跑工程正确性全集，容差只能来自工程现有测试，禁止自创业务容差；任一目标用例失败即转精度调试，不进入性能环节。
2. 每轮门禁（Step 7）：每轮单变量修改后先跑完整精度再跑同口径性能——精度失败该轮直接回滚，该轮性能数字不得出现在报告结论中。
3. 结论门禁（Step 8）：delta_report.py 在脚本层强制校验前后 precision=pass，否则拒绝生成对比报告；result_saver.py 落盘 mismatch/maxAbsErr/maxRelErr 误差指纹，前后精度数字必须同量级。

校验链本身必须先审计再采信（实测踩坑：某 FIA 案例官方校验脚本对 NaN 位型退化为"全一致即 pass"，NaN 掩码是死代码）：

• 跑通校验脚本 ≠ 校验有效。确认：golden 生成链在当前软件栈真实可用（标杆 aclnn 拒绝该场景即断裂）；比对方式是位级还是数值转换（数值转换可能先销毁 NaN/Inf）；掩码/跳过分支是否真的执行。
• 校验链断裂或退化时：停止性能调优，改用语义仿真 golden（CPU 参考）并保持工程原容差，补充无异常位型的辅助数据集，处置过程写入报告。

改动按语义风险分级审查：

• 语义保持类（分核/blockDim、双缓冲、搬运合并、同步删减、调度顺序）：数学结果不变，常规精度门禁即可。
• 语义变更类（算法替换、近似公式、dtype/量化精度转换、归约顺序改变可能影响浮点结合律）：除门禁外，必须在候选方案中写明精度风险与额外验证用例（边界值、异常位型、大 shape），收益结论需标注精度代价。

Step 3：建立不可变基线

固定源码基线、软件栈、SoC、设备、频率、shape、dtype、format、物理 padding、TilingKey、blockDim、warmup、repeat 和计时方法。优先使用设备 event 或工程正式 benchmark；msOpProf 只用于结构诊断，短 kernel 不得用 profiler 插桩耗时替代 event 基线。

计时口径的边界情形（实测补充）：

• 工程无 benchmark/计时设施（cann-samples 样例普遍如此）：允许在产物目录的工程副本上给 host 侧补 aclrtEvent 计时 harness（warmup/repeat 固定），属测量设施而非优化变量；模板与归类规则见 benchmark-harness.md。原始基线源码必须先备份。
• 长 kernel（ms 级）且工程无 event 设施：允许用 msopprof OpBasicInfo Task Duration 作基线（result_saver --timing-method msprof_task_duration），报告中注明口径；工程 README 的参考数字同口径时才可对比。
• 短 kernel（µs 级单 launch）：launch 开销会掩盖 kernel 级收益（wall 口径下任何优化都不可归因）。用 stream 内多 launch 单 event 对计时；仍不可归因时以 profiler device duration 作辅助口径并注明，不要按 wall 噪声判回滚或宣称收益。
• 冷跑/稳态差异：首跑含初始化与缓存预热（实测可差 25%），固定 warmup 剔除；稳态 ≥3 次取中位，run 间噪声 >5% 时小于噪声幅度的收益不得归因。
• 工程自带 profiler 对比框架（如 torch_npu op_summary 的 torch vs triton benchmark）：视为工程正式 benchmark，按本节优先级采信。

python3 <skill_dir>/scripts/result_saver.py \
  --op <operator-or-case> --variant baseline --mode board --soc <a2|a3|a5> \
  --precision pass --kernel-avg-us <value> --timing-method event \
  --baseline-kind <system|self_before_after> --cann-version <version> \
  --repo-commit <baseline-commit> --device-id <id> \
  --shape <shape-or-case-id> --dtype <dtype> --format <format> \
  --tiling-key <key> --block-dim <n> --warmup <n> --repeat <n> \
  --msprof-dir <path-or-empty> --output tuning_artifacts/<operator>/baseline

没有等价系统算子时使用 baseline_kind=self_before_after，不得构造 builtin 对比。

Step 4：采集 msOpProf

读取 profile/profile-msopprof.md，先探测本机帮助再生成命令。必须：

1. 用 kernel 过滤条件和结果清单确认所有目标 kernel 均被采集。
2. 保存 OpBasicInfo.csv 与本机支持的基础 Pipe/Memory 指标。
3. A2/A3 可优先评估 MemoryDetail、TimelineDetail；A5 可优先评估 PipeTimeline、PcSampling，但只使用本机帮助明确支持的项。
4. MC2/多 rank 仅在本机明确支持时使用 msprof op，否则回退完整 msprof。
5. 保存采集命令、输出目录、kernel 过滤、launch 次数和失败项。

Step 5：证据化判定 Bound

按已有数据读取：

• 核间长尾或 block 利用率：diagnose/diagnose-occupancy.md
• Compute/Memory/Latency 分类：diagnose/diagnose-roofline.md
• L2/GM 流量与工作集：diagnose/diagnose-l2-cache.md
• MTE/Vector/Cube/Scalar/FixPipe 流水：diagnose/diagnose-pipeline.md

有 PipeUtilization.csv 时执行：

python3 <skill_dir>/scripts/bound_analyzer.py \
  --csv <PipeUtilization.csv> \
  --output tuning_artifacts/<operator>/baseline/bound_report.json

脚本只做初筛。最终结论必须包含 Bound、CSV/字段/数值/核或 kernel、排除其他 Bound 的反证、置信度和缺失指标。高 busy 必须结合 active bandwidth，区分真实带宽上限与小包、非对齐或同步频繁。

Step 6：选择同型优化机制

读取 优化技术索引，只加载命中文档：

证据或 DSL	资料
切核不均、尾块长尾	Ascend C Tiling
MTE2/MTE3、小包或非对齐	数据搬运
L2/UB/L1/L0 复用不足	存储层次
Pipe 气泡或同步过密	流水优化
Scalar/API 热点	API 使用
多步融合 Vector、MemBase 已尽、VF 样例	RegBase / SIMD VF
CATLASS / Triton / TileLang	CATLASS / Triton / TileLang
PyPTO / SHMEM	PyPTO / SHMEM

再从 案例路由 选择最多三个同 DSL、同数据流、同 SoC 能力且同 Bound 的案例。公开数字只作为来源事实。候选项必须写明对应证据、修改位置、预期指标、架构限制、精度风险、验证命令和回滚条件。

Step 7：单变量迭代

每轮只实施一个可归因机制或一组不可拆分的原子修改：保存 diff，执行最小改动，以同一命令重编译，先跑完整精度再跑同口径 benchmark，最后补采受影响指标。精度失败、稳定劣化或收益落在噪声内时回滚并记录原因。

负优化是正常且可处置的（21 案例实测中出现率约 1/3）：机制与瓶颈不匹配、容量/深度联动假设被破坏（tiling 改动后 fault）、同步协议改动死锁、或收益被噪声淹没，都会导致某轮变慢或不可用。本流程用四道机制兜住：

1. 可预防：候选机制必须有当前算子的 profiling 证据（不是案例数字外推）；单变量保证可归因；原子修改对先二分验证单路径可用性。
2. 可检测：同口径对比（delta_report.py 硬口径校验）+ 噪声带判定（±5% 内不算收益也不算劣化，需加采样轮次）+ 冷跑/稳态区分。
3. 可回滚：每轮有 diff、基线源码有备份，回滚后必须复跑确认回到基线水平（验证回滚干净），再进入下一轮。
4. 可交代：回滚的轮次与原因写进最终报告的"保留和回滚的修改"——负优化记录本身是瓶颈证据的组成部分，不是失败。

不硬编码固定轮数。达到目标、候选耗尽、连续两轮无稳定收益或证据显示接近硬件上限时停止。

Step 8：前后对比

按 Step 3 的完整参数保存优化后结果；repo_commit 使用优化后的提交，其余比较口径必须一致。然后执行：

python3 <skill_dir>/scripts/delta_report.py \
  --baseline <baseline-result.json> --after <optimized-result.json> \
  --output tuning_artifacts/<operator>/comparison.md

最终报告包含环境与源码指纹、完整用例表、baseline/optimized/Delta/speedup/精度、Bound 与关键指标变化、保留和回滚的修改、官方文档或 PR 及适用边界。无法采集或未验证的内容标记为 partial 或 dry_run。

失败分流与禁止事项

编译、运行或采集失败时读取 troubleshooting.md。禁止：精度失败时继续调优；未审计校验链即采信精度结论；跨软件栈、SoC、频率、shape、dtype 或 TilingKey 比较；无证据宣称 Bound；照搬案例加速比；一次修改多个无关机制；劣化或噪声内收益的轮次不回滚、不记录；只展示最快样本或提升 case；批量加载 references/cases/。