# Triton · Programming models and profiling 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** (A style of parallel program in which the code the compiler works with operates on blocks of data at a time, rather than on one thread's worth of data at a time.) 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** (Scheduling iteration blocks statically, based on the control- and data-flow structure of the program, rather than at the level of individual threads.) 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. > Figure. Scalar program versus blocked program. On the left, CUDA's framing: you write one thread's worth of scalar arithmetic, the hardware replicates it across the block, and each thread computes one output element. On the right, Triton's: every statement already operates on a whole tile, and the coalescing, swizzling, shared-memory, and thread-mapping decisions a CUDA programmer makes by hand move into the compiler.Illustrative numbers 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: > Figure. add_kernel, from the Triton vector addition tutorial 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](https://learn-kernels.com/chapters/reading#triton) for tutorials and the original paper. Source Blocked-program framing and the matmul comparison from the [Triton programming guide](https://triton-lang.org/main/programming-guide/chapter-1/introduction.html), chapter 1. Project description from the [triton-lang/triton](https://github.com/triton-lang/triton) README. The add\_kernel listing, mask, and grid mechanics from the [vector addition tutorial](https://triton-lang.org/main/getting-started/tutorials/01-vector-add.html), with comments trimmed. Triton-to-TileIR from [NVIDIA’s CUDA Tile IR backend for Triton](https://developer.nvidia.com/blog/advancing-gpu-programming-with-the-cuda-tile-ir-backend-for-openai-triton/) post.