Triton 算子静态代码检视（Ascend NPU）
检视原则
- Ascend 特有约束优先：聚焦硬件差异，不检查 Triton 通用知识
- Mask 零容错：Ascend 对越界访问零容忍
严重性分级
| 级别 | 含义 | 典型问题 |
|---|---|---|
| P0 | 必定崩溃或错误结果 | Mask 遗漏、核类型错配、Atomic 循环死锁 |
| P1 | 高概率精度/功能问题 | 归约未升精度、Softmax 未减 max |
| P2 | 性能/可维护性 | 冗余访存、BLOCK 未对齐 |
参考资源加载
| 阶段 | 加载 | 不要加载 |
|---|---|---|
| Phase 1: Host 侧 | ascend-triton-api-constraints.md | dtype-matrix |
| Phase 2: Device 侧 | ascend-api-dtype-matrix.md | test-patterns |
| 逐项核对 | code-review-checklist.md | — |
| 参考官方实现 | ascend-test-patterns.md | — |
检视工作流
Phase 1: Host 侧
| 检查项 | 识别方式 | 级别 |
|---|---|---|
| 硬编码核数 | grid = (20,) 等字面量 | P0 |
| 核类型错配 | 含 tl.dot 的 kernel 用了 num_vectorcore | P0 |
BLOCK_SIZE 非 tl.constexpr | 声明检查 | P1 |
| 矩阵运算 BLOCK 非 16 倍数 | 数值检查 | P2 |
核类型：含 tl.dot → AI Core (num_aicore)；逐元素/归约/激活 → Vector Core (num_vectorcore)
Phase 2: Device 侧
Mask 完整性（P0）：所有 tl.load/tl.store 必须有 mask= 或使用 make_block_ptr。
数据类型（P0-P1）：tl.dot 输入仅支持 int8/fp16/fp32/bf16；dot_scaled 不支持；permute/trans 不支持 int64。
精度处理（P1）：FP16/BF16 归约前 .to(tl.float32)；Softmax 必须减 max。
代码模式：
- ❌
for...: tl.atomic_cas/or/xor/and/xchg(...)— 可能死锁（P0） - ❌ 多核 kernel 中使用
tl.atomic_add返回值（P0） - ❌ kernel 内
import numpy（P0） - ⚠️
tensor[i].item()在 Host 热路径 — 触发 CPU-NPU 同步（P2）
Phase 3: 性能隐患（P2）
- 同一 ptr 多次
tl.load→ 冗余 GM 访问 tl.arange(0, N) * stride（stride > 1）→ 非连续访存pid直接映射 block 无循环 → 负载不均- kernel 内对辅助张量（cos/sin 等）使用 broadcast stride 计算偏移（
b * stride0 + n * stride1 +...）→ 建议改为 host 侧 expand+contiguous，统一row * D + col偏移
反模式清单（NEVER）
Host
- ❌ 硬编码核数 / 矩阵乘法用
num_vectorcore/ BLOCK_SIZE 非tl.constexpr
Device
- ❌
tl.load/tl.store无 mask /tl.dot输入 int32/int16/int64 /dot_scaled - ❌
atomic_or/xor/and/xchg/cas在 for 循环内 / kernel 内第三方库 - ❌ FP16/BF16 归约不升 FP32 / Softmax 不减 max
输出
按 code-review-report-template.md 格式输出报告。