Triton programming model

Programs, not threads

A Triton kernel is a function over programs. The launch grid says how many programs exist, and tl.program_id(0) says which one is running:

@triton.jit
def copy_kernel(x_ptr, y_ptr, n, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)   # a block of lane indices
    mask = offs < n
    tl.store(y_ptr + offs, tl.load(x_ptr + offs, mask=mask, other=0.0), mask=mask)

tl.arange(0, BLOCK) is not a loop. It is a block of BLOCK lanes, and every operation below applies to all of them at once. There is no threadIdx.x, no __syncthreads, and no shared memory in the source. num_warps decides how many threads cooperate on one program, and the compiler decides how lanes map onto them.

Masks replace index arithmetic

From CUDA you would write if (i < n) around the load. Here the predicate becomes a mask of the same shape as the block, passed to tl.load and tl.store, with other supplying the value a masked lane holds. That value matters: a max over other=0.0 returns zero for an all-negative row, and the bug looks like a numerical quirk rather than a wrong kernel.

What the compiler owns

You choose the tile shape and the masks. Triton chooses the distribution of the block across threads and registers, whether a dot stages operands through shared memory, how loads are vectorised, and how the loop is pipelined. That trade is the point of the language: the tile is the unit of thought, and warp-level detail stays the compiler’s problem until you read the TTGIR.

Things that break

A block size that is not a power of two fails at compile time, because Triton requires it. The tail is where the interesting failures live: a missing mask either faults or reads padding that happens to be allocated and returns plausible numbers, which is how a wrong kernel passes a test on a lucky allocator. These snippets follow the Triton 3.9 documentation; confirm tl.load and the default num_warps against triton.__version__ before copying them, because both have moved.