Triton: Program Tiles Instead of Manual Thread Mapping

Core Concept

A Triton program instance processes a logical data tile. The author describes tile shapes, addresses, masks, and reductions while the compiler maps work to warps and threads.

Why It Affects Performance

Triton reduces manual scheduling effort and is well suited to fusion and row reductions, but coalescing, tile size, register pressure, and numerical behavior still require design.

Example or Pseudocode

@triton.jit
def saxpy(x, y, out, a, n: tl.constexpr, BLOCK: tl.constexpr):
    offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n
    tl.store(out + offs, a*tl.load(x+offs, mask=mask) + tl.load(y+offs, mask=mask), mask=mask)

Profiling and Performance Judgment

Warm up JIT compilation, benchmark with Triton helpers or CUDA Events, autotune BLOCK and num_warps, and inspect generated behavior with Nsight.

Five-Minute Exercise

Does BLOCK=1024, num_warps=4 mean 1,024 CUDA threads?