Triton Operator Performance Eval

Triton 算子性能评估（Ascend NPU）

核心认知

性能数据信仰：只相信测量到的数据，不相信自己的直觉和假设。评估前必须用 Prof 定位真正的性能热区。

⚠️ 唯一可信的性能采集方式：msprof 和 msprof op

通过任何非 msprof / msprof op 方式（包括但不限于 Python time.time()、torch.npu.Event 计时、triton.testing.do_bench、自定义计时装饰器等）采集的性能数据绝不可接受，绝不可用于性能评估及优化决策。这些方式精度不足、无法排除系统调度和 JIT 编译等干扰因素，得出的数据不具备任何参考价值。所有性能分析必须且只能基于 msprof（函数级）或 msprof op（算子级）的输出数据。

评估目标：

• 识别性能瓶颈类型（Memory-Bound vs Compute-Bound）
• 量化硬件利用率
• 对比不同实现的性能差异
• 验证优化效果

性能评估工作流

函数级性能采集（首选）

使用 msprof 进行函数级性能分析：

msprof --application="python my_script.py" --output=./profiling_result

适用场景：

• 对比多个 PyTorch 算子 vs 融合 Triton 算子的性能
• 分析函数级别的性能瓶颈
• 生成可视化性能报告
• 全链路性能分析（Host + Device）

详细用法和示例：加载 msprof-function-level.md

算子级性能采集（深度分析）

使用 msprof op 进行算子级深度分析：

msprof op --kernel-name={jit_kernel_name} {application}

适用场景：

• 分析单个 Triton kernel 的硬件利用率
• 诊断 Cube/Vector 单元性能
• 分析 UB 缓存和内存访问模式
• 定位 Bank Conflict 等硬件问题

详细用法和示例：加载 msprof-op-level.md

性能数据分析

关键指标：

• 瓶颈类型判断（Memory-Bound vs Compute-Bound）
• 内存带宽利用率
• 计算单元利用率（Cube/Vector）
• UB 冲突分析

详细分析方法：加载 performance-data-analysis.md

性能总结输出

完成性能评估后，必须按以下模板输出结构化性能总结：

## 性能评估总结

### 基本信息
| 项目 | 值 |
|------|-----|
| 算子名称 | {kernel_name} |
| 输入规模 | {shape, dtype} |
| 测试硬件 | {Ascend 型号} |
| 测量方法 | {msprof / msprof op} |

### 性能指标
| 指标 | 值 | 参考值 | 利用率 |
|------|-----|--------|--------|
| 执行耗时 | {X} us | - | - |
| 内存带宽 | {X} GB/s | {理论峰值} GB/s | {X}% |
| Cube 利用率 | - | - | {X}% |
| Vector 利用率 | - | - | {X}% |
| L2 Cache 命中率 | - | - | {X}% |
| Bank Conflict 比例 | - | - | {X}% |

### 瓶颈诊断
- **瓶颈类型**：{Memory-Bound / Compute-Bound}
- **判断依据**：{基于 Arithmetic Intensity 和硬件利用率数据的分析}
- **关键证据**：{引用具体 CSV 数据}

### 性能问题清单
| 优先级 | 问题 | 证据 | 优化方向 |
|--------|------|------|----------|
| P0 | {最关键问题} | {数据来源} | {具体建议} |
| P1 | ... | ... | ... |

### 优化建议
1. {最高优先级优化建议及预期收益}
2. ...

输出原则：

• 所有结论必须有 Profiling 数据支撑，不做主观猜测
• 利用率低于 30% 的指标标记为 重点关注
• 对比场景需同时列出 baseline 和 optimized 结果

参考资源加载指南

MANDATORY - 按需加载：根据任务类型加载对应的参考文档

性能评估检查清单

基础检查

• 是否使用 msprof 进行性能采集？
• 是否进行了预热（避免首次编译开销）？
• 是否多次测量取统计值？
• 是否同步了 NPU 设备（torch.npu.synchronize()）？

Profiler 检查

• 是否指定了正确的 --application 或 --kernel-name？
• 是否选择了合适的 --aic-metrics？
• 是否分析了所有关键性能指标？
• 是否识别了瓶颈类型（Memory-Bound vs Compute-Bound）？

性能指标检查

• 内存带宽利用率是否合理？
• 计算单元利用率是否合理？
• 是否存在高 Bank Conflict？
• L2 Cache 命中率是否合理？

反模式清单（NEVER）

• NEVER 使用任何非 msprof / msprof op 的方式进行计时或性能评估（包括 time.time()、torch.npu.Event、triton.testing.do_bench、自定义计时器等）——这些方式采集的数据精度不足、无法排除系统调度和 JIT 编译干扰，绝不可接受，绝不可用于任何性能评估及优化决策
• NEVER 在性能测试中不进行预热（首次执行包含编译开销）
• NEVER 只测试一次就得出结论（需要多次测量取统计值）
• NEVER 在性能测试中包含打印或日志（I/O 会严重影响结果）
• NEVER 忘记同步 NPU 设备（torch.npu.synchronize()）
• NEVER 在不同硬件环境下对比性能结果
• NEVER 混淆 msprof 和 msprof op 命令（前者函数级全局分析，后者算子级深度分析）
• NEVER 在没有 Profiling 数据支撑时给出优化建议

常见陷阱与注意事项

性能问题与优化方向

瓶颈类型与优化策略

常见性能问题诊断

优化方向速查

Memory-Bound 算子优化路径：

1. 检查访存模式 → 确保连续访存（Memory.csv 带宽利用率）
2. 减少数据搬运 → 算子融合减少 GM 读写次数
3. 提高数据复用 → 优化 Tiling 策略使数据在 UB/L1 中多次使用
4. 消除 Bank Conflict → 调整对齐方式（ResourceConflictRatio.csv）

Compute-Bound 算子优化路径：

1. 命中 Cube 单元 → BLOCK 维度设为 16 倍数（ArithmeticUtilization.csv）
2. 减少类型转换 → 避免不必要的 upcast/downcast
3. 流水线优化 → 检查 Pipe 利用率，平衡计算和搬运（PipeUtilization.csv）
4. 向量化 → 确保 Vector 操作充分利用 SIMD 宽度

参考资源

msprof vs msprof op 命令对比

这两个命令是完全不同的分析层级，选择错误会导致无效分析：

选择决策：

• 需要知道"哪个算子最慢" → 用 msprof
• 需要知道"这个 kernel 为什么慢" → 用 msprof op
• 完整优化流程：先 msprof 定位热点，再 msprof op 深度分析

参考文档

官方资源

• Triton-Ascend 官方文档
• Triton-Ascend 教程
• Triton 官方文档
• msprof 官方文档