What problem does it solve? Triton-Ascend kernels written with Cube/Vector affinity often fail verification due to subtle bugs: a single UB buffer overwritten across loop iterations (WAW), bl.alloc shapes that are runtime tensors instead of constexpr, and pure matmul kernels wrongly written in affinity form. This Skill documents the root causes and concrete fixes for these recurring precision and compilation failures. ## Core Features & Use Cases - WAW Buffer Overwrite Fix: Explains why a single c_ub buffer reused across outer-loop iterations yields only the last written tile, and provides the paired sync_block_set/wait pattern that lets cube and vector scopes safely reuse the buffer. - Constexpr Shape Enforcement: Diagnoses the get_buffer_ty TypeError from bl.alloc and shows two fixes—annotating parameters with tl.constexpr or decoupling runtime dimensions behind a BLOCK constant with masked loops. - Affinity Misuse Correction: Identifies when a no-postprocessing GEMM should not use al.scope/fixpipe at all, and provides the native Triton tl.dot + tl.make_block_ptr rewrite. - Use Case: A kernel engineer sees err_cnt mismatches on an Atlas A5 matmul kernel using fixpipe to UB; this Skill pinpoints the WAW root cause and supplies the corrected synchronization code. ## Quick Start Ask the agent to diagnose why my Triton-Ascend affinity kernel fails verification and apply the appropriate fix from this guide.
