triton language

Triton is OpenAI's open-source language and compiler for writing GPU kernels in Python. It sits in the gap between calling a black-box library like cuBLAS and hand-writing CUDA C++: you describe what one program instance does to a *block* of data, and the compiler handles the thread-level parallelism, memory coalescing, shared-memory staging, and instruction scheduling that a CUDA programmer would otherwise manage by hand. (This is the *Triton language*, not NVIDIA's separately-named Triton Inference Server, which is an unrelated model-serving product.)\n\n**Triton's core idea is to raise the unit of programming from the thread to the block.** In CUDA you write code from the point of view of a single thread and reason explicitly about `threadIdx`, warps, and `__shared__` memory. In Triton you write code from the point of view of one *program* in a launch grid, and every operation acts on a whole tile: `tl.load` pulls a `BLOCK_SIZE`-wide slice through a pointer and a boolean mask, arithmetic runs elementwise over the tile, and `tl.store` writes it back. The compiler decides how to spread that tile across threads and warps, so the same source runs well across different block sizes and hardware generations.\n\n**You address memory with pointers and masks instead of thread indices.** A Triton kernel receives raw pointers plus tensor strides, computes a vector of offsets like `pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)`, and loads with a mask that guards the ragged tail of a non-divisible dimension. This is what lets Triton generate coalesced, vectorized loads automatically: because the access pattern is expressed as arithmetic over a contiguous tile, the compiler can prove it is regular, emit wide aligned transactions, and manage the shared-memory buffers for reductions and matmul accumulation without the programmer writing a single `__syncthreads()`.\n\n**Autotuning is a first-class part of the workflow.** Kernel performance on a GPU is dominated by a few meta-parameters: the tile shape (`BLOCK_M/N/K`), how many warps execute one program (`num_warps`), and how many pipeline stages overlap global loads with compute (`num_stages`). The `@triton.autotune` decorator sweeps a list of these configurations, benchmarks them for each new input shape, and caches the winner. This replaces the CUDA ritual of hand-templating over launch bounds, and it is why a few dozen lines of Triton can match a vendor kernel that took an expert weeks to tune.\n\n**Under the hood Triton is an MLIR-based compiler, not a source-to-source translator.** A `@triton.jit` function is traced into Triton IR, lowered to TritonGPU IR (a dialect that carries tile layouts and warp-level information), then to LLVM IR and finally to PTX/SASS for NVIDIA, with AMD and other backends maturing. The middle stages are where the real work happens: software pipelining of load-then-compute, allocation of shared memory, layout conversions between tensor-core-friendly and register-friendly forms, and vectorization. This is the same machinery that PyTorch's `torch.compile` targets: its Inductor backend *emits Triton* for the fused GPU kernels it generates, so Triton is increasingly the substrate that ordinary PyTorch code lowers down to.\n\n**Triton earns its keep on fusion, not on replacing BLAS.** The kernels people reach for Triton to write are the ones no library ships: a fused softmax, a matmul with a custom epilogue, layer-norm-plus-residual in one pass, or the tiled online-softmax at the heart of FlashAttention. Fusing these into a single kernel keeps intermediates in registers and shared memory instead of round-tripping through HBM, which is exactly where memory-bound models spend their time. For a plain dense GEMM the vendor library is usually still the right call; Triton wins when the shape is unusual, the epilogue is custom, or several operations can be melted together.\n\n| Approach | You program at the level of | Shared memory & sync | Iteration speed | Best when |\n|---|---|---|---|---|\n| cuBLAS / cuDNN | a library call | vendor-managed | instant | standard dense GEMM / conv |\n| **Triton** | a **block / tile** | **compiler-managed** | fast (Python + autotune) | fused and custom kernels |\n| CUDA C++ | a single **thread** | you, by hand | slow (recompile, hand-tune) | exotic patterns, the last 5% |\n\n```svg\n \n```\n\nRead Triton through a *what-does-one-block-do* lens rather than a *what-does-one-thread-do* lens: you are describing tile-level intent and letting an MLIR compiler synthesize the thread choreography, which is why a short, hackable kernel can land within a few percent of a hand-tuned vendor library and why it has become the compilation target underneath PyTorch itself.