triton-op-coding

Provide the operator name, the task description containing the reference Model class, and the target architecture such as A2. The Skill loads operator-type-specific references and outputs a complete Python file with @triton.jit kernels and a ModelNew class matching the original Model's interface.

GPU Triton kernels can only serve as structural references for signatures, grid patterns, and data flow. GPU-specific parameters like num_warps, num_stages, and num_ctas have no effect on NPU and must be removed in favor of Ascend configurations such as num_cores.

Common causes include missing fp32 accumulation in reductions, incorrect mask handling on non-divisible tails, or mismatched random weights. For operators with nn.Linear or nn.Conv2d, ModelNew must call torch.manual_seed(0) first and recreate modules in the exact original order.

All core computation must live inside @triton.jit kernels. The forward method may only allocate buffers with torch.empty, perform shape operations like view or permute, query metadata like shape and dtype, and launch kernels—torch.matmul, tensor methods, and arithmetic operators are prohibited.

When the project template's Layer 1 constraints conflict with the sketch architecture, the Layer 1 hard rules take priority. The generator redesigns the code structure accordingly, preserves non-conflicting sketch details like tile sizes, and annotates each deviation in code comments.

Use an interleaved loop where a fixed grid equal to the core count strides over blocks, or a continuous block method mapping one grid to multiple data chunks. Core counts are queried once in __init__ via the driver device properties API, never inside forward.