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\nTriton'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\nYou 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\nAutotuning 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\nUnder 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\nTriton 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.