Reading production kernel code

01 · The checklist

Six questions, in order, every time

Production kernels look intimidating — hundreds of lines, template metaprogramming in the CUDA ones, autotune decorators and specialization keys in the Triton ones. But they all answer the same six questions, and reading them in order turns strangeness into structure:

02 · A distilled kernel

RMSNorm, the way it appears in vLLM — minus the production armor

RMSNorm normalizes each row by its root-mean-square, then applies a learned per-channel weight. It runs once or twice per layer in every modern LLM — a hot path worth reading. Below is distilled from vLLM's Triton RMSNorm: structure and line inventory preserved, production details (dtype dispatch, caching allocators, stride handling for non-contiguous tensors) trimmed:

@triton.jit
def rms_norm(x_ptr, w_ptr, y_ptr, n_cols, eps, BLOCK: tl.constexpr):
    row = tl.program_id(0)
    offs = tl.arange(0, BLOCK)
    mask = offs < n_cols

    x = tl.load(x_ptr + row * n_cols + offs, mask=mask,
                other=0.0).to(tl.float32)
    var = tl.sum(x * x, axis=0) / n_cols
    rstd = 1.0 / tl.sqrt(var + eps)

    w = tl.load(w_ptr + offs, mask=mask)
    tl.store(y_ptr + row * n_cols + offs, x * rstd * w, mask=mask)

Run the checklist. Output tile: one row per program — question 1 answered by line 3. Input traffic: row-contiguous loads, coalesced by construction (day 20); the weight vector is read by every program and lives in L2 (day 10). Accumulation: tl.sum(x*x) — day 16's reduction, compiler- lowered. Staging: the row sits in registers for its whole lifetime — no shared memory needed at one row per program. Tuning: BLOCK is a constexpr padded to a power of two; in the real file it's wrapped in autotune configs (day 22). The profile: two rows in flight worth of traffic per program — read + write — over trivial math: memory-bound, near the bandwidth ceiling, exactly as day 11 predicts. One line earns a second look: the .to(tl.float32) upcast — summing squares in FP16 loses precision long before the result looks wrong, so production kernels widen before reducing. Day 25's accumulate-wide rule, wearing a different dtype.

03 · Where to read

A reading list, ordered by difficulty

Start where the grammar is simplest and climb. The Triton tutorials (fused attention, layer norm, matmul with autotuning) are annotated and gentle. Then vLLM and SGLang: their Triton kernels — RMSNorm, rotary embeddings, sampling, paged attention (day 26, now readable) — are production-hot but structurally familiar, and the CUDA C++ side (csrc/) shows the same kernels at lower altitude. The FlashAttention repository comes after: day 24 gave you the idea; the source shows the epilogue handling, causal masking, and GQA bookkeeping that the idea has to survive in production. CUTLASS is graduate school — go when a GEMM article stops scaring you.

Reading order within any kernel: find the launch site first (grid shape, block shape — question 1), then the tile mapping, then the loads. Never read a kernel top-down like an essay; read it like a map, from the territory you recognize outward.

04 · What you're looking for

Patterns repeat until you can predict them

After a handful of kernels you'll notice the field is mostly remixes: tile-map, load, maybe stage, reduce or accumulate, store. When you can hear a kernel's name — rotary_embedding, fused_add_rms_norm, paged_attention_v1 — and sketch its likely structure before opening the file, day 01's promise is paid in full: production kernels are no longer foreign text. They're your grammar, spoken with an accent. Day 28 starts the capstone by going the other direction — not reading speed, but making it.

Next: day 28 — the Capstone begins. Profile and optimize: taking day 16's reduction from ~20% of bandwidth to ~88%, one measured change at a time.