external-gitcode-ascend-triton-operator-code-review

Triton Operator Static Code Review (Ascend NPU)

Review Principles

• Ascend-specific constraints take priority: Focus on hardware differences, do not check general Triton knowledge
• Zero tolerance for Mask: Ascend has zero tolerance for out-of-bounds access

Severity Classification

Level	Meaning	Typical Issues
P0	Definite crash or incorrect result	Missing Mask, core type mismatch, Atomic loop deadlock
P1	High-probability precision/functional issues	Reduction without precision promotion, Softmax without max subtraction
P2	Performance/maintainability	Redundant memory access, BLOCK not aligned

Reference Resource Loading

Phase	Load	Do not load
Phase 1: Host side	ascend-triton-api-constraints.md	dtype-matrix
Phase 2: Device side	ascend-api-dtype-matrix.md	test-patterns
Item-by-item verification	code-review-checklist.md	—
Reference official implementation	ascend-test-patterns.md	—
Confirm API signature/parameters	triton-api-reference.md	—

Load trigger: When it is necessary to confirm the signature, parameters, or available data types of a certain tl/libdevice API, read in full triton-api-reference.md。

Review Workflow

Phase 1: Host side

Check item	Identification method	Level
Hardcoded core count	grid = (20,) and other literals	P0
Core type mismatch	Contains tl.dot kernel uses num_vectorcore	P0
Matrix multiplication degraded to element-wise	No tl.dot, implementing matmul/GEMV with Vector Core element-wise multiply-add	P0
BLOCK_SIZE not tl.constexpr	Declaration check	P1
Matrix operation BLOCK not a multiple of 16	Value check	P2

Core type: Contains tl.dot → AI Core (num_aicore); element-wise/reduction/activation → Vector Core (num_vectorcore). Note: matrix multiplication operators (including GEMV) must use tl.dot + AI Core, Cube Core throughput is much higher than Vector Core.

Phase 2: Device side

Mask completeness (P0): All tl.load/tl.store must have mask= or use make_block_ptr。

Data types (P0-P1)：tl.dot input supports only int8/fp16/fp32/bf16;dot_scaled conditionally supported (bf16/fp16 lhs/rhs, int8 scale, fp32 out);permute/trans does not support int64.

Precision handling (P1): FP16/BF16 before reduction .to(tl.float32); Softmax must subtract max.

Single-pass check for reduction operators (P1): For operators such as GroupNorm/LayerNorm/RMSNorm, check whether the same data is loaded multiple times (first compute statistics, then load for normalization). When D ≤ UB, it should be single-pass: after one load, tl.sum(x, 1) complete all computations within UB. Double-pass can be 5x slower.

Control flow (P0): Triton IR uses MLIR structured control flow (scf.for/scf.while), does not support early exit.

• ❌ for/while Inside the loop return — Compile error: "Cannot have return statements inside while or for" (including return in subfunctions)
• ❌ for Loop-internal break — Compile error: "unsupported AST node type: Break"

Tensor indexing (P0): Triton tensors do not support Python-style [] subscript operations.

• ❌ tensor[i] = val — AssertionError
• ❌ val = tensor[i] — AssertionError
• ❌ tensor[i:j] slicing — Compile error

Code patterns：

• ❌ for ... : tl.atomic_cas/or/xor/and/xchg(...) — possible deadlock (P0)
• ❌ Using tl.atomic_add return values in multi-core kernels (P0)
• ❌ In-kernel import numpy（P0）
• ⚠️ tensor[i].item() on Host hot path — triggers CPU-NPU synchronization (P2)
• ⚠️ Reduction-type operators have multiple for ... : tl.load(x_ptr + ...) loops — double-pass anti-pattern (P1)
• ⚠️ Post-dot operations do not use tl.parallel(bind_sub_block=True) — tl.dot followed by element-wise operations but no parallel sharding (P2)
• ⚠️ Integer type conversion overflow — tl.cast to int8/int16 without specifying overflow_mode（P1）

Phase 3: Performance pitfalls (P2)

• Same ptr multiple tl.load → redundant GM access
• tl.arange(0, N) * stride(stride > 1) → non-contiguous memory access
• pid directly maps block without loop → load imbalance
• In-kernel use of broadcast stride to compute offsets for auxiliary tensors (cos/sin, etc.) (b * stride0 + n * stride1 + ...) → recommend changing to host-side expand+contiguous, unify row * D + col offsets
• Reduction-type operators weight/bias not pre-expanded, in-kernel use of weight_ptr + ch_idx gather → recommend host-side expansion to contiguous layout

Anti-pattern list (NEVER)

Host

• ❌ Hardcoded core count / matrix multiplication uses num_vectorcore / BLOCK_SIZE not tl.constexpr
• ❌ Matrix multiplication (including GEMV) implemented with Vector Core element-wise multiply-add — must use tl.dot + AI Core

Device

• ❌ tl.load/tl.store without mask / tl.dot input int32/int16/int64 / dot_scaled
• ❌ atomic_or/xor/and/xchg/cas inside for loop / third-party libraries in kernel
• ❌ FP16/BF16 reduction without promoting to FP32 / Softmax without subtracting max
• ❌ Integer truncation cast without specifying overflow_mode
• ❌ for/while inside loop return / break — Triton structured control flow does not support early exit (including return in subfunctions)
• ❌ tensor[i] indexing operations (read/write/slice) — use tl.where / tl.gather / tl.extract_slice instead

Output

Output the report in code-review-report-template.md format.