LAB / GPU PROGRAMMING

DAY 19 / 30 · TRITON & MODERN KERNEL AUTHORING · ~12 MIN READ

Why Triton exists

The abstraction gap between hand-indexed CUDA and block-level programming.

01 · The tax of days 13–18

Everything you hand-rolled, every kernel, forever

Look back over the Writing CUDA phase with slightly colder eyes. Thread indexing and bounds guards in every kernel (days 13–14). Coalescing by careful index arithmetic (day 07). Shared-memory staging with paired barriers, padded against bank conflicts (days 08, 15). Block sizes argued about, then tuned (day 14). The reductions, the shuffles, the partial passes (day 16). Each was worth learning — day 27 will depend on all of it — but notice a pattern: most of this work is formulaic and repeats in every kernel you'll ever write. The genuinely novel part of each kernel is usually ten lines in the middle.

In 2019, OpenAI's Philippe Tillet made the obvious-but-radical bet: if the per-lane work is formulaic, the compiler should do it. Triton is that bet, compiled: a Python-embedded language where you write block-level programs and the compiler writes the thread-level ones.

02 · Blocks, not threads

The same kernel, one abstraction level up

In Triton you never name a thread. Your function describes what happens to one block of data — a tile, in day 15's vocabulary — and the runtime launches it once per block, like CUDA launching a kernel once per… well, once per program. Vectors of offsets replace per-thread indices; masks replace guards; tl.load and tl.store move whole tiles at a time. It is day 15's tiled thinking with the phase-loop bookkeeping deleted.

Pick a concern below and watch the same small kernel pay its tax in CUDA versus Triton — the highlighted lines are who owns that concern:

WHO OWNS WHAT · SAME KERNEL, TWO LEVELS OF ABSTRACTION

CUDA · YOU OWN IT

__global__ void scale_add(const float* x, float* out,
                          int n, float a) {
  // 1. thread indexing — yours, every kernel
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = gridDim.x * blockDim.x;
  for (; i < n; i += stride) {
    // 2. bounds masking — yours, every kernel
    float v = 0.0f;
    if (i < n) v = x[i];
    // 3. coalescing — yours (did x[i] map to lanes?)
    // 4. staging — yours (tile[ty][tx] = …; __syncthreads();)
    out[i] = a * v;
  }
}  // 5. tuning — yours (<<<blocks, threads>>>)

TRITON · COMPILER OWNS IT

@triton.jit
def scale_add(x_ptr, out_ptr, n, a, BLOCK: tl.constexpr):
  # 1. indexing is block-level — one line
  pid = tl.program_id(0)
  offs = pid * BLOCK + tl.arange(0, BLOCK)
  # 2. masking — one expression
  mask = offs < n
  # 3. coalescing — compiler lowers lanes for you
  # 4. staging — compiler stages if (and only if) it helps
  v = tl.load(x_ptr + offs, mask=mask, other=0.0)
  tl.store(out_ptr + offs, a * v, mask=mask)
  # 5. tuning — @triton.autotune searches BLOCK, num_warps

CUDA: per-thread global id plus a grid-stride loop (day 14). Triton: one program id and a block of offsets — the compiler cuts it into lanes.

simplified: both programs are complete and equivalent; the CUDA side elides the actual staging code for brevity — which is rather the point.

03 · What you trade away

No lanes, no layouts, no mercy for exotic tricks

The price is control. You cannot express warp-level shuffle tricks (day 16's last 32 steps are the compiler's business now), cannot dictate shared-memory layouts or padding, cannot hand-schedule tensor-core fragments or re-order instructions. If the compiler's cost model makes a bad call, you adjust the knobs it exposes — block shapes, warp counts, pipeline stages (day 22) — not the assembly underneath.

For most kernels this is a spectacular trade: fused elementwise chains, normalizations, softmaxes, rotary embeddings, small-to-medium GEMMs, attention variants. It's why PyTorch's torch.compile emits Triton, and why vLLM and SGLang ship Triton kernels for their hottest paths. The frontier exceptions — cuBLAS-class GEMM, cutting-edge attention — still belong to hand-written CUDA and CUTLASS, where specialists fight for the last 20%. Day 24's FlashAttention sits close enough to the frontier that its authors wrote CUDA; the fused softmax of day 21 does not.

04 · Choosing your tool

The decision is about distance from the frontier

A working rule: start in Triton, drop to CUDA when the profiler proves the compiler is the bottleneck. Triton's iteration speed — Python, no build step, autotuning — makes the first 80% of performance arrive in an afternoon; the last 20% is where hand-tuned CUDA buys its keep, if the kernel is important enough to need it. And because day 27's reading list is mostly vLLM and SGLang, learning to read Triton first is simply learning the dialect you'll encounter most.

CHEATSHEET · DAY 19
  • 01CUDA's per-lane work is formulaic — so let the compiler own it. Indexing, masking, coalescing, staging, and tuning repeat in every kernel. Triton moves all five below your line of sight.
  • 02You write one program per block; the compiler writes the lanes. Block-shaped offsets and masks replace thread indices and guards. Day 15's tile thinking, minus the bookkeeping.
  • 03The price is lane-level control. No warp shuffles, no shared-memory layouts, no hand-scheduled tensor cores. Frontier kernels still drop to CUDA.
  • 04Start Triton; drop to CUDA when profiled, not before. torch.compile, vLLM, and SGLang all emit or ship Triton — it's the dialect of modern open-source kernels.

Next: day 20 — Triton's programming model properly: three primitives, masks, block-shaped pointers, and what the compiler actually does with your tile.

get in touch