# Triton Operator Performance Eval

> 评估 Ascend NPU 上 Triton 算子的性能表现。当用户需要分析算子性能瓶颈、使用 msprof/msprof op 进行算子性能采集与对比、诊断 Memory-Bound/Compute-Bound 瓶颈、测量硬件利用率指标、生成性能评估报告时使用。

- Skill: `small-cat/triton-operator-performance-eval` (Agent Skill)
- Install (CLI): `npx skillmds@latest add small-cat/triton-operator-performance-eval`
- Raw SKILL.md: https://api.skillmd.com/api/skills/small-cat/triton-operator-performance-eval/raw
- Safety review: pending
- Works with: Claude Code, Claude.ai, OpenAI Codex
- Author: small-cat (https://skillmd.com/u/small-cat)
- Updated: 2026-09-22
- Page: https://skillmd.com/skills/small-cat/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 进行函数级性能分析：**

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

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

**详细用法和示例**：加载 [`msprof-function-level.md`](references/msprof-function-level.md)

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

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

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

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

**详细用法和示例**：加载 [`msprof-op-level.md`](references/msprof-op-level.md)

### 性能数据分析

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

**详细分析方法**：加载 [`performance-data-analysis.md`](references/performance-data-analysis.md)

### 性能总结输出

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

```markdown
## 性能评估总结

### 基本信息
| 项目 | 值 |
|------|-----|
| 算子名称 | {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-function-level.md` | `msprof-op-level.md` |
| 算子级硬件分析 | `msprof-op-level.md`, `performance-data-analysis.md` | `msprof-function-level.md` |
| 性能瓶颈诊断 | `performance-data-analysis.md` | - |
| 理解硬件术语 | `ascend-terminology.md` | - |
| 完整性能优化流程 | 所有 references | - |

## 性能评估检查清单

### 基础检查
- [ ] 是否使用 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 数据支撑时给出优化建议

## 常见陷阱与注意事项

| 陷阱 | 表现 | 正确做法 |
|------|------|----------|
| 用 `msprof op` 做算子对比 | 只能看到单个 kernel，无法对比 | 算子对比用 `msprof`，深度分析用 `msprof op` |
| `--kernel-name` 拼写错误 | `msprof op` 静默完成但无数据 | 确认 kernel 名称与 Triton 函数定义一致 |
| 未区分首次编译和稳态性能 | 首次运行耗时异常高 | 至少 5 次预热后再采集 |
| 小规模输入测性能 | 启动开销占比过大，结论无参考价值 | 使用生产规模输入进行评估 |
| 忽略 dtype 对性能影响 | FP16 和 FP32 性能差异显著 | 固定 dtype 进行对比，分别评估 |

## 性能问题与优化方向

### 瓶颈类型与优化策略

| 瓶颈类型 | 判断条件 | 核心优化方向 |
|----------|----------|-------------|
| **Memory-Bound** | AI < 硬件平衡点；带宽利用率高、计算利用率低 | 减少数据搬运量、提高数据复用、优化访存模式 |
| **Compute-Bound** | AI > 硬件平衡点；计算利用率高、带宽利用率低 | 优化计算指令效率、提高 Cube/Vector 利用率 |
| **Latency-Bound** | 带宽和计算利用率均低 | 增大并行度（Grid Size）、减少同步开销 |

### 常见性能问题诊断

| 问题 | 症状 | 诊断数据源 | 解决方向 |
|------|------|-----------|----------|
| UB 溢出 | 编译错误/运行时 OOM | 检查 BLOCK_SIZE 配置 | 减小 BLOCK_SIZE 或核内再分块 |
| Cube 未命中 | 性能仅 10% 理论值 | ArithmeticUtilization.csv | 强制 BLOCK_M/N/K=16 倍数 |
| 精度损失 | FP16 结果偏差大 | 对比 PyTorch 结果 | 累加器用 FP32 |
| 非连续访存 | 带宽仅 20% 利用率 | Memory.csv | 调整数据布局为连续 |
| 低并行度 | AI Core 利用率低 | PipeUtilization.csv | 增大 Grid Size |
| 高 Bank Conflict | 资源冲突率 > 10% | ResourceConflictRatio.csv | 调整数据块大小和对齐方式 |
| L2 Cache 命中率低 | 频繁 GM 访问 | L2Cache.csv | 优化 Tiling 策略，提高数据局部性 |

### 优化方向速查

**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`（函数级） | `msprof op`（算子级） |
|------|-------------------|---------------------|
| **命令格式** | `msprof --application="python x.py"` | `msprof op --kernel-name=K python x.py` |
| **分析粒度** | 整个应用的所有算子 | 指定的单个 kernel |
| **核心输出** | op_summary.csv、timeline_trace.json、report.html | ArithmeticUtilization.csv、Memory.csv、PipeUtilization.csv 等 |
| **提供信息** | 各算子耗时排名、Host/Device 全链路时间线 | 硬件利用率、内存带宽、Bank Conflict 等微架构指标 |
| **典型用途** | 对比 PyTorch vs Triton 算子整体性能 | 诊断单个 kernel 的硬件瓶颈 |
| **必需参数** | `--application` | `--kernel-name` |

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

### 参考文档

| 文档 | 内容 | 关联命令 |
|------|------|---------|
| [`msprof-function-level.md`](references/msprof-function-level.md) | 函数级性能采集用法和输出分析 | `msprof` |
| [`msprof-op-level.md`](references/msprof-op-level.md) | 算子级深度分析用法和硬件指标 | `msprof op` |
| [`performance-data-analysis.md`](references/performance-data-analysis.md) | `msprof op` 输出 CSV 的详细分析方法 | `msprof op` |
| [`profiling-tools.md`](references/profiling-tools.md) | 性能分析工具链总览和工作流 | 两者均涉及 |
| [`ascend-terminology.md`](references/ascend-terminology.md) | Ascend 硬件术语和架构概念 | - |

### 官方资源
- [Triton-Ascend 官方文档](https://gitcode.com/Ascend/triton-ascend/blob/main/docs/zh)
- [Triton-Ascend 教程](https://gitcode.com/Ascend/triton-ascend/tree/main/python/tutorials)
- [Triton 官方文档](https://triton-lang.org/main/index.html)
- [msprof 官方文档](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta1/devaids/Profiling/atlasprofiling_16_0010.html)