Ascend/triton-operator-code-review

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 | — | | 确认 API 签名/参数 | triton-api-reference.md | — |

加载触发：需要确认某个 tl/libdevice API 的签名、参数或可用数据类型时，完整阅读 triton-api-reference.md。

检视工作流

Phase 1: Host 侧

| 检查项 | 识别方式 | 级别 | |--------|---------|------| | 硬编码核数 | grid = (20,) 等字面量 | P0 | | 核类型错配 | 含 tl.dot 的 kernel 用了 num_vectorcore | P0 | | 矩阵乘法退化为逐元素 | 无 tl.dot，用 Vector Core 逐元素乘加实现 matmul/GEMV | P0 | | BLOCK_SIZE 非 tl.constexpr | 声明检查 | P1 | | 矩阵运算 BLOCK 非 16 倍数 | 数值检查 | P2 |

核类型：含 tl.dot → AI Core (num_aicore)；逐元素/归约/激活 → Vector Core (num_vectorcore)。注意：矩阵乘法类算子（含 GEMV）必须用 tl.dot + AI Core，Cube Core 吞吐远高于 Vector Core。

Phase 2: Device 侧

Mask 完整性（P0）：所有 tl.load/tl.store 必须有 mask= 或使用 make_block_ptr。

数据类型（P0-P1）：tl.dot 输入仅支持 int8/fp16/fp32/bf16；dot_scaled 有条件支持（bf16/fp16 lhs/rhs, int8 scale, fp32 out）；permute/trans 不支持 int64。

精度处理（P1）：FP16/BF16 归约前 .to(tl.float32)；Softmax 必须减 max。

归约类算子单 pass 检查（P1）：GroupNorm/LayerNorm/RMSNorm 等算子，检查是否对同一数据多次 load（先算统计量，再 load 做归一化）。当 D ≤ UB 时应单 pass：一次 load 后 tl.sum(x, 1) 在 UB 内完成全部计算。双 pass 可慢 5x。

控制流（P0）：Triton IR 使用 MLIR 结构化控制流（scf.for/scf.while），不支持 early exit。

• ❌ for/while 循环内 return — 编译错误："Cannot have return statements inside while or for"（包括子函数中的 return）
• ❌ for 循环内 break — 编译错误："unsupported AST node type: Break"

Tensor 索引（P0）：Triton tensor 不支持 Python 风格 [] 下标操作。

• ❌ tensor[i] = val — AssertionError
• ❌ val = tensor[i] — AssertionError
• ❌ tensor[i:j] 切片 — 编译错误

代码模式：

• ❌ 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）
• ⚠️ 归约类算子存在多个 for ... : tl.load(x_ptr + ...) 循环 — 双 pass 反模式（P1）
• ⚠️ Post-dot 操作未用 tl.parallel(bind_sub_block=True) — tl.dot 后有逐元素操作但无并行分片（P2）
• ⚠️ 整数类型转换溢出 — tl.cast 到 int8/int16 未指定 overflow_mode（P1）

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 偏移
• 归约类算子 weight/bias 未预展开，kernel 内用 weight_ptr + ch_idx gather → 建议 host 侧展开为连续布局

反模式清单（NEVER）

Host

• ❌ 硬编码核数 / 矩阵乘法用 num_vectorcore / BLOCK_SIZE 非 tl.constexpr
• ❌ 矩阵乘法（含 GEMV）用 Vector Core 逐元素乘加实现——必须用 tl.dot + AI Core

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
• ❌ 整数截断 cast 未指定 overflow_mode
• ❌ for/while 循环内 return / break — Triton 结构化控制流不支持 early exit（含子函数中的 return）
• ❌ tensor[i] 索引操作（读取/赋值/切片）— 用 tl.where / tl.gather / tl.extract_slice 替代

输出

按 code-review-report-template.md 格式输出报告。