Building a GPU Compiler from Scratch

A step-by-step walkthrough shows how to lower PyTorch graphs to CUDA by progressively transforming through multiple intermediate representations. It lays out six stages - from Torch IR that captures framework ops, through a unified Tensor IR that makes indexing/reductions explicit, to Loop IR that fuses scalar loops, Tile IR that separates computation from schedule, Kernel IR that spells hardware work, and finally emitted CUDA for nvcc. Concrete examples (FP16 matmul and RMSNorm) and command-line snippets for the Emmy compiler illustrate how transposes, broadcasts, and reductions are represented with index maps and explicit sums, how fusion replaces intermediate memory with in-loop values, and how a computation tree with nested Folds maps to a schedule describing work distribution and reduction placement.

The walkthrough then exposes scheduling as a codec of discrete knobs and demonstrates tuning tradeoffs with two concrete rows for a 128×128 FP16 matmul (choices for warps per tile, register-tile shape, reduction chunking, and staging depth). It shows how different knobs change data reuse versus occupancy and why measurement is required to pick winners. Lowering a chosen schedule to CUDA produces helpers and inline asm (cp.async, mbarrier, ldmatrix) that implement shared-memory staging and tensor-core fragments. The central claim is that a multi-level IR stack plus an explicit schedule representation makes exploration, fusion, and hardware mapping tractable for building a practical GPU compiler.