Part IV — Back end

Code generation

Lesson 08 chose a schedule for the laid-out graph: tile sizes, loop order, what to stage into SRAM, where to vectorize and unroll. That is a complete recipe for how the loops should run — and not one byte of it executes. A schedule is a plan written in the compiler's own notation; the GPU runs instructions, and there is no instruction in a schedule. This lesson is where the plan becomes machine code: codegen, the lowering of a scheduled IR onto a real backend. The fork in the road is which backend — emit low-level LLVM IR → PTX → SASS and let LLVM do register allocation and instruction selection, or emit a tile DSL like Triton and let its compiler map tiles onto warps and threads. Modern stacks took the second road, and this lesson is the argument for why.

1 · What codegen produces: device code + host launch code

Codegen does not emit "a kernel." It emits two coupled things, and forgetting the second one is the classic beginner's blind spot. The device code is the kernel itself — the program each GPU thread runs, expressed eventually as SASS (the GPU's actual machine instructions). The host code is the CPU-side glue that launches it: it allocates or reuses output buffers, computes the launch grid (how many thread blocks, of what shape), packs the kernel arguments (pointers, strides, scalar params), and issues the launch on a CUDA stream. The schedule from lesson 08 fixes both: the tile size sets the grid (grid = ⌈M/BLOCK_M⌉ × ⌈N/BLOCK_N⌉), and the loop body becomes the device kernel.

Where eager PyTorch issued one library kernel per op with the launch handled by ATen (cross-link gpu_kernels · 18 (PyTorch to kernel)), the compiler must now generate both halves itself — because the fused kernels from lesson 05 do not exist in any library; only the compiler knows their shape. That is the whole point of getting here: bespoke kernels demand bespoke codegen.

2 · Path A — lower to LLVM IR → NVVM → PTX → SASS

The classic path treats the GPU like any other compiler target: lower the scheduled loop nest to LLVM IR — a low-level, typed, SSA instruction set with explicit loads, stores, arithmetic, and branches, the same IR Clang uses for C++ — then hand it to a target backend. For NVIDIA that backend is NVVM (NVIDIA's LLVM-based compiler), which emits PTX (Parallel Thread Execution: a stable, virtual assembly / ISA that is forward-compatible across GPU generations), which the driver's ptxas finally compiles down to SASS (the actual, architecture-specific machine code for one GPU, e.g. sm_90 for Hopper). XLA, TVM's LLVM backend, and Triton's own bottom all end here.

What LLVM gives you for free is enormous and is exactly why people reach for it: register allocation (deciding which values live in the GPU's limited registers vs spill to local memory), instruction selection (matching IR patterns to the best target opcodes), constant propagation, peephole optimization, and a mature backend per ISA. You describe the computation scalar-by-scalar; LLVM turns it into a tight register schedule.

What LLVM does not give you is the hard part of tensor code. LLVM reasons at the level of scalars and machine vectors; it has no native concept of a tile, of shared-memory staging, of warp-level tensor-core matmul, or of the cooperative dance where 256 threads load a 128×128 block, multiply it on the tensor cores, and write it back. To get a fast matmul out of the LLVM path you must hand-write all of that thread choreography — the shared-memory double-buffering, the mma fragment layouts, the bank-conflict-free swizzles — directly, in PTX or in CUDA-with-intrinsics. The IR is low; the abstraction it offers (a scalar machine with vectors) is the wrong shape for a tensor program. That is the gap the next path closes.

scheduled loop nest (lesson 08)
        │  lower scalars, loads/stores, branches
        ▼
   LLVM IR ──► regalloc + instruction selection  (LLVM gives you this)
        │      but: no tiles, no SRAM staging, no warp-MMA  (you hand-write these)
        ▼
   NVVM ──► PTX (virtual ISA, forward-compatible)
        │      ptxas
        ▼
   SASS (sm_90 machine code) ──► the SM executes it

3 · Path B — emit a tile DSL (Triton) and let it map threads

The second path raises the codegen target from a scalar machine to a tile machine. Instead of lowering all the way to LLVM yourself, you emit Triton — a Python-embedded kernel language whose unit of work is the tile, not the thread. You write the program for one block: which output tile it owns, which input tiles to tl.load into SRAM, the block-level matmuls and reductions on that hot data, the tl.store back. Triton's own compiler (built on MLIR) then does the part LLVM could not: it maps your tiles onto warps and threads, lays out shared memory, picks tensor-core mma instructions, schedules the loads to hide latency, and finally lowers that to LLVM → PTX → SASS. You handed off the thread choreography. The deeper treatment of the tile model and how Triton lowers is in gpu_kernels · 22 (writing Triton).

This is why Inductor — the codegen backend of torch.compile — generates Triton rather than raw CUDA/PTX (cross-link gpu_kernels · 23 (torch.compile)). The compiler's job after fusion (lesson 05) is to emit a kernel per fusion group, and there are thousands of distinct groups across real models. Emitting hand-tuned PTX for each is intractable; emitting Triton is a few dozen lines of templated codegen. The trade is explicit:

The verdict the field reached: for the long tail of fused pointwise/reduction kernels — the bulk of what a compiler emits — Triton's portability and tiny codegen surface dominate, and the few percent of peak you leave on the table is recovered by fusing more aggressively and by autotuning (lesson 10). Raw PTX/SASS is reserved for the handful of kernels where peak is worth a human's month: the flagship GEMM, FlashAttention (cs336/07), the cuDNN/cuBLAS-class library primitives. Everything else, the compiler generates as Triton.

4 · Worked: the fused bias + gelu, generated two ways

Take the fused bias + gelu kernel from lesson 05 — the elementwise epilogue that, fused, reads the activation from HBM once, adds the bias, applies GELU, writes once. Here is what Inductor-style codegen emits for it: a Triton kernel. Note the level — tiles, masks, block-wide math; no thread, no register named.

@triton.jit
def fused_bias_gelu(X, B, Y, n_elements, n_features, BLOCK: tl.constexpr):
    pid  = tl.program_id(0)                       # this block owns one tile
    offs = pid * BLOCK + tl.arange(0, BLOCK)      # the tile's element indices
    mask = offs < n_elements                      # guard the ragged last tile
    x = tl.load(X + offs, mask=mask)              # one HBM read of the tile
    b = tl.load(B + offs % n_features, mask=mask) # broadcast bias along features
    h = x + b                                      # bias add, in registers
    # tanh-approx GELU, all on hot data — never written to HBM:
    g = 0.5 * h * (1.0 + tl.math.tanh(0.79788456 * (h + 0.044715 * h * h * h)))
    tl.store(Y + offs, g, mask=mask)             # one HBM write of the tile
# host: grid = (triton.cdiv(n_elements, BLOCK),)   <-- launch code, generated too

A handful of readable lines of math plus a host launch line. The compiler picks BLOCK (and num_warps) — that is what lesson 10 tunes. Now contrast the shape of the same kernel down the LLVM/PTX path. You do not write tiles; you write the scalar loop, and the generated PTX is per-thread, verbose, and ISA-specific:

// sketch of the same bias+gelu as generated PTX (per-thread, abridged)
.visible .entry fused_bias_gelu(.param .u64 X, .param .u64 B, .param .u64 Y, .param .u32 n) {
    mov.u32    %r1, %ctaid.x;          // block id
    mov.u32    %r2, %ntid.x;           // block dim
    mov.u32    %r3, %tid.x;            // thread id
    mad.lo.s32 %r4, %r1, %r2, %r3;     // global index  = block*dim + thread
    setp.ge.s32 %p1, %r4, %n;          // bounds guard
    @%p1 bra   DONE;
    // ... compute byte offset, ld.global.f32 x, ld.global.f32 b ...
    add.f32    %f3, %f1, %f2;          // bias add
    // ... ~20 more PTX instrs: h*h*h, tanh poly, 0.5*h*(1+...) ...
    st.global.f32 [%rd_y], %f_out;     // store
DONE: ret;
}

Two facts jump out. First, length: the Triton version is ~6 lines of intent; the PTX is dozens of per-thread opcodes, and that ratio explodes for anything with a matmul or a reduction, where the PTX must spell out shared-memory staging and mma fragments by hand. Second, lock-in: the PTX is NVIDIA-only and you would write a separate file for AMD; the Triton compiles to either. For a single elementwise kernel the gap is annoying. For the thousands of fusion groups a compiler emits, it is the difference between a maintainable backend and an impossible one — which is precisely why Inductor emits the top snippet, not the bottom.

5 · Cache the artifact: compile once, run a million times

Codegen and the subsequent LLVM/ptxas (or Triton) compile are not free — emitting and compiling one kernel costs from a few tens of milliseconds (a simple Triton kernel) to seconds (a heavily-autotuned GEMM under max-autotune). If that cost were paid on every forward pass, the compiler would be a catastrophe. It is not, because of the one fact that makes the whole enterprise pay off: an ML workload runs the same shapes a million times. So codegen compiles each kernel once and stores the compiled artifact in a code cache, keyed by everything that would change the generated code — the source, the input shapes and dtypes, the schedule config, the target architecture. Every later call with a matching key is a cache hit: the kernel is already a binary, and you just launch it.

This is why torch.compile's first step is slow and the rest fly (cross-link gpu_kernels · 23): the first call is a compile storm, the steady state is pure cache hits. Two consequences matter in production. Cold start: a latency-sensitive service must warm up with representative shapes before serving, or the first user pays the compile. Cache key: anything that changes the key — a new shape, a new dtype — is a miss and a recompile, which is exactly the wall that dynamic shapes (lesson 11) crash into. For now, hold the win: codegen pays its cost once and the cache turns it into a rounding error.

6 · Drive it: pick the codegen target

The widget below makes the codegen-target trade a knob. For the fused bias + gelu (and, toggle on, a tiled matmul to see the gap widen), pick the backend — LLVM → PTX/SASS or Triton — and the tile/num_warps config. Watch four readouts: lines of generated code (the codegen-surface cost), portability (which hardware the artifact runs on), est. peak-% reachable (how close to hand-tuned), and compile time. The knob that breaks: pick PTX for the matmul and watch lines-of-code and compile time explode while portability collapses to one ISA — the few extra points of peak come at a cost no compiler can pay across thousands of kernels.

Slide it and the verdict is stable: Triton sits near the floor on lines-of-code and compile time and runs on both vendors at ~90% of peak; PTX climbs to the last few points of peak only by paying a 10–40× line-count penalty, a slower compile, and a one-ISA lock-in. For the one GEMM that runs in every layer, a team pays that. For the thousands of fused kernels behind it, the compiler emits Triton and moves on.

Failure modes & checklist

Checkpoint

Where this points next

Codegen closes the gap from plan to machine code: a scheduled loop nest becomes a device kernel plus host launch, emitted as Triton for the long tail and reserved-PTX for the flagships, then cached so its compile cost amortizes to nothing. But notice what we waved past. Codegen can faithfully emit any schedule you hand it — which silently assumes you already chose the right one. Lesson 08 picked a schedule by hand for one matmul; a real model has thousands of kernels, each with a configuration space of tile sizes (powers of two from 16 to 256 on each axis), num_warps (1–16), num_stages (pipeline depth), and unroll factors — a combinatorial space whose best point depends on the exact shape and the exact GPU. No human picks BLOCK=256, num_warps=4 per kernel per shape, and the widget's knobs you just slid are precisely the ones that need choosing. The only way out is to search the space — benchmark candidates, or predict them with a cost model — and keep the winner. That is lesson 10, Autotuning & cost models.

Interview prompts

• Codegen produces more than "a kernel." What are the two parts and why does the second matter? (§1 — the device kernel (lowered to SASS) and the host launch code that computes the grid from the tile size, packs args, and issues the launch; without the host half nothing runs, and for fused kernels neither half exists in any library so the compiler must generate both.)
• Walk the classic lowering path and say what each level is. (§2 — scheduled IR → LLVM IR (typed SSA, scalar-level) → NVVM → PTX (virtual, forward-compatible ISA) → SASS (real per-architecture machine code via ptxas). LLVM does regalloc + instruction selection.)
• What does LLVM give you, and why is it the wrong abstraction for tensor code? (§2 — it gives register allocation, instruction selection, peephole opts; but it reasons in scalars/vectors with no native tile, SRAM staging, or warp-MMA concept, so you must hand-write the thread choreography to get tensor-core speed.)
• Why did Inductor choose to generate Triton instead of raw CUDA/PTX? (§3–4 — it must emit a kernel per fusion group, thousands of them; Triton is a tile-level DSL where its MLIR compiler maps threads, so codegen is a few templated lines per group and the artifact is portable (NVIDIA+AMD), trading ~5–15% peak for a maintainable backend.)
• Distinguish PTX from SASS. (§2 — PTX is a stable, forward-compatible virtual ISA emitted by NVVM; SASS is the real, architecture-specific machine code that ptxas produces from PTX and the SM actually executes. Cache SASS per arch, or you JIT-recompile.)
• Codegen + compile cost tens of ms to seconds per kernel. Why isn't that fatal? (§5 — ML runs the same shapes a million times, so each kernel is compiled once and stored in a code cache keyed by source/shapes/dtypes/config/arch; later calls are cache hits that just launch the cached binary, amortizing compile to ~nothing.)
• What is the cache key, and which later lesson does it set up? (§5 — everything that would change generated code: source, input shapes, dtypes, schedule config, target arch; a new shape is a miss → recompile, which is exactly the dynamic-shapes wall in lesson 11.)