Triton

Triton is, in its own words, “a language and compiler for writing highly efficient custom Deep-Learning primitives,” aimed at giving programmers “higher productivity than CUDA” and “higher flexibility than other existing DSLs.” It grew out of a MAPL 2019 paper by Tillet, Kung, and Cox, and the premise stated in its own documentation is that blocked programs can produce compute kernels competitive with hand-written CUDA, while staying far more flexible than prior compiler frameworks like Halide or TVM.

The distinction Triton draws against CUDA is precise. CUDA, in its own framing, is “Scalar Program, Blocked Threads”: a programmer writes one thread’s worth of scalar arithmetic, and the hardware replicates it across a block of threads, each computing one output element of a matrix multiply. Triton inverts this to “Blocked Program, Scalar Threads”: the program itself loops over tiles of the output, stepping by block sizes MB, NB, and KB, and every statement inside that loop already operates on whole blocks of A, B, and the accumulator at once. There is no per-thread code to write, because the block is the unit the language exposes.

That inversion only pays off because of what the compiler does with it. Triton relies on block-level data-flow analysis to turn a blocked program into a scheduled one, and that analysis is what drives automatic memory coalescing, thread swizzling, prefetching, vectorization, tensor-core-aware instruction selection, shared-memory allocation and synchronization, and asynchronous copy scheduling. Those are exactly the decisions a CUDA programmer makes by hand at the thread-block level described in chapter 1; Triton moves them into the compiler and asks the programmer to reason about tiles instead.

Triton’s own first tutorial makes the style concrete with the same computation chapter 1 used to introduce CUDA: vector addition. Where vecAdd gave every thread exactly one element, the Triton kernel is written per program, and each program handles BLOCK_SIZE elements at once. The tutorial’s comments call out the two moves that define the model: the kernel asks tl.program_id which of the parallel program instances it is, and offsets “is a list of pointers,” a whole block of addresses that tl.load reads in one statement:

@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements,
               BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    output = x + y
    tl.store(output_ptr + offsets, output, mask=mask)


def add(x: torch.Tensor, y: torch.Tensor):
    output = torch.empty_like(x)
    n_elements = output.numel()
    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )
    add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
    return output

Nothing in the kernel names a thread. The mask guards the loads and stores against out-of-bounds accesses when the vector length is not a multiple of the block size, which is the blocked model’s replacement for chapter 1’s per-thread bounds check. The launch side is what the tutorial calls an SPMD grid, “analogous to CUDA launch grids”: here a 1D grid of cdiv(n_elements, BLOCK_SIZE) program instances, with each torch.Tensor argument implicitly converted into a pointer to its first element. Everything below the program, the warps, the shared memory, the vector widths, is the compiler’s problem.

That same tile abstraction is why Triton and CUTLASS keep meeting in the same place. NVIDIA has since built a backend, Triton-to-TileIR, that lets Triton programs compile directly to CUDA Tile IR (covered next) instead of PTX, preserving Triton’s tile semantics rather than lowering them to per-thread SIMT instructions first. See further reading for tutorials and the original paper.