Warmup Guide — GPU Compilers: NVCC, PTX, SASS, and Codegen

Zero-to-expert primer for Phase 03. Assumes Phase 02 (you've compiled CUDA with nvcc) and basic compiler vocabulary is NOT assumed — we build it. By the end you can read PTX/SASS, explain the JIT/AOT production tradeoffs, and write a small compiler that emits legal PTX.

Table of Contents

Chapter 1: Why a Platform Leader Must Understand the Compiler Layer

The JD's opening mandate — "a platform that operates independently of the underlying hardware" — is a sentence about compilers. Hardware independence means some layer translates a stable program representation into each vendor's machine code. Whoever controls that layer controls the platform's performance, its portability claims, and its cost structure. Concretely:

  • When vLLM gains 30% from a new attention kernel, that's the compiler/kernel layer.
  • When a startup accelerator "supports PyTorch" but delivers 25% of paper FLOPs, the gap is compiler maturity (Phase 01 Ch. 10's question 5).
  • When your service cold-starts for minutes on a new GPU generation, that's JIT compilation you didn't plan for (Ch. 7 — a real incident pattern).
  • When NVIDIA's dominance is described as "the CUDA moat," the load-bearing piece is PTX + 18 years of kernels that still run (Ch. 4, 10).

You don't need to be a compiler engineer. You need to never be lied to about this layer — by vendors, by candidates, or by your own roadmap.

Chapter 2: Compilers From Zero — the Five Classic Stages

For readers who've never built one (you will, in Lab 02):

  1. Lexing: characters → tokens. c[i] = a[i] * 2.0 becomes IDENT(c) LBRACK IDENT(i) RBRACK ASSIGN IDENT(a) ... NUM(2.0).
  2. Parsing: tokens → AST (a tree mirroring the grammar). Recursive descent — one function per grammar rule — is how you'll write it; precedence climbing handles * binding tighter than +.
  3. Semantic analysis / type checking: annotate the AST with types (f32, i32), reject nonsense, decide where conversions go.
  4. IR + optimization: a flat, instruction-like intermediate representation — in real compilers, SSA form (every value assigned exactly once; makes data-flow analysis trivial). Passes rewrite it: constant folding, common subexpression elimination (CSE), dead-code elimination (DCE).
  5. Codegen: IR → target instructions + register allocation + the target's calling convention/ABI.

The deep insight you'll internalize by building one: a compiler is just a chain of small, testable tree/list transforms. None of the stages are magic; all of the difficulty is in the target's details — which is why the rest of this guide is about the target.

Chapter 3: The NVCC Pipeline, For Real

nvcc is an orchestrator. Run nvcc -dryrun x.cu once in your life and read the plan it prints. The flow:

x.cu ── cudafe++ ──► host C++ ─────────────► g++/clang ──► host .o ─┐
   └──► device C++ ──► cicc (NVVM/LLVM) ──► x.ptx                   ├─► executable
                       ptxas ─────────────► x.cubin (SASS)          │
                       fatbinary ─────────► fatbin (ptx+cubins) ────┘
  • cudafe++ splits host/device code (the __global__/__device__ markers).
  • cicc is NVIDIA's LLVM-based device frontend: C++ → NVVM IR (an LLVM IR dialect) → optimizations → PTX.
  • ptxas is the real backend: PTX → SASS (the architecture-specific ISA), doing register allocation, scheduling, and the occupancy-relevant decisions (-Xptxas=-v from Phase 02 is asking ptxas to talk).
  • fatbinary bundles PTX (for forward-compat JIT) + one or more cubins (AOT) into the executable; the runtime picks at load time (Ch. 7).

The same split exists in every GPU stack: a frontend to a portable IR, then a per-architecture backend. Hold the shape; vendors only change the names.

Chapter 4: PTX — The Virtual ISA That Is CUDA's Moat

What it is: PTX (Parallel Thread eXecution) is a virtual instruction set — assembly for a GPU that doesn't physically exist, designed to be stable across hardware generations.

Properties that make it virtual (and that your Lab 02 compiler will emit):

  • Infinite typed registers: .reg .f32 %f<N> — declare as many as you like; ptxas does the real allocation. (This is why register pressure problems are invisible in PTX and only appear in SASS — Lab 01 step 3.)
  • Explicit state spaces: every memory access names its space — .param (kernel arguments), .global, .shared, .local, .const. The cvta.to.global instruction converts a generic address to a space-specific one. PTX makes Phase 01's memory hierarchy syntactically visible.
  • Explicit types on every instruction: add.f32, mad.lo.s32, ld.global.f32 — no inference at the ISA level.
  • Predication instead of (most) branches: any instruction can be guarded: @%p1 bra DONE; or even @%p st.global... — Phase 01's active-mask machinery, visible in text.
  • Special registers: %tid.x, %ctaid.x, %ntid.x — threadIdx, blockIdx, blockDim. The Phase 02 index formula compiles to exactly: mad.lo.s32 %r4, %ctaid.x, %ntid.x, %tid.x.

Why "moat": NVIDIA promises PTX forward compatibility — PTX compiled in 2010 JITs onto an H100 today. That promise is what lets two decades of CUDA software accumulate without ever being rewritten, and it's the single hardest thing for any competitor to replicate, because it's not a chip feature — it's an ecosystem contract (Ch. 10).

Chapter 5: Reading PTX — A Guided Tour of saxpy

y[i] = a * x[i] + y[i] — the kernel you'll annotate in Lab 01. Abridged, with the narration you should be able to give:

.visible .entry saxpy(            // __global__ visible to the driver
    .param .f32 saxpy_param_0,    // a        — args live in .param space
    .param .u64 saxpy_param_1,    // x ptr
    .param .u64 saxpy_param_2,    // y ptr
    .param .u32 saxpy_param_3     // n
){
    ld.param.f32  %f1, [saxpy_param_0];
    ld.param.u64  %rd1, [saxpy_param_1];
    cvta.to.global.u64 %rd2, %rd1;     // generic -> global address space
    mov.u32  %r1, %ctaid.x;            // blockIdx.x
    mov.u32  %r2, %ntid.x;             // blockDim.x
    mov.u32  %r3, %tid.x;              // threadIdx.x
    mad.lo.s32 %r4, %r1, %r2, %r3;     // i = blockIdx*blockDim + threadIdx
    setp.ge.s32 %p1, %r4, %r5;         // p1 = (i >= n)
    @%p1 bra DONE;                     // the tail guard, as a predicated branch
    mul.wide.s32 %rd3, %r4, 4;         // byte offset = i * sizeof(float)
    add.s64 %rd4, %rd2, %rd3;          // &x[i]
    ld.global.f32 %f2, [%rd4];         // x[i]
    fma.rn.f32 %f4, %f1, %f2, %f3;     // a*x[i] + y[i]  — contracted to FMA
    st.global.f32 [%rd5], %f4;         // y[i] = ...
DONE:
    ret;
}

Five things to notice (they're the Lab 01 annotation rubric):

  1. Arguments arrive via ld.param — the kernel ABI. Your Lab 02 compiler emits this exact prologue.
  2. cvta.to.global — the state-space conversion; pointers from the host are "generic" until proven global.
  3. The index computation is a single mad — multiply-add, integer flavor.
  4. The bounds check is a predicate register + predicated branch — and for tiny bodies the compiler skips the branch entirely and predicates the store (if-conversion, Ch. 8).
  5. fma.rn.f32 — the compiler contracted a*x + y into one fused multiply-add, round-to-nearest. (This is why GPU and CPU results differ in the last ulp — Phase 02's tolerance-based verification, explained.)

Chapter 6: SASS — Where the Machine Stops Pretending

SASS is the real, per-architecture ISA (cuobjdump --dump-sass). What changes versus PTX:

  • Real registers: R0–R254 + predicates; allocation is done; spills appear as LDL/STL (local loads/stores) — the smoking gun Lab 01 step 3 hunts after -maxrregcount 16.
  • Control information embedded: modern SASS encodes stall counts, dependency barriers, and reuse hints chosen statically by ptxas — the GPU does no out-of-order discovery (Phase 01 Ch. 1's deleted machinery, visible as compiler homework instead).
  • Architecture-specific instructions: HMMA/QGMMA (Tensor Core ops), LDGSTS (the async copy behind cp.async), wider loads LDG.E.128.
  • No stability promise: SASS encodings change every generation — which is why PTX exists, and why "ship PTX in the fatbin" is the forward-compat story.

You don't write SASS; you read it for three signals: spills (LDL), whether Tensor Cores were used (HMMA), and what the inner loop actually became (unrolled? predicated? vectorized loads?). That reading level is reachable in an afternoon — Lab 01 gets you there.

Chapter 7: JIT vs AOT, Fatbins, and the Cold-Start Incident

The compatibility model — a production operations topic disguised as compiler trivia:

  • -arch=compute_80 -code=sm_80,sm_90,compute_80 → fatbin contains: SASS for sm_80, SASS for sm_90, and PTX (compute_80).
  • At load, the runtime picks: exact SASS match → run it. No match but PTX present → the driver JIT-compiles PTX → SASS on first load, caching in ~/.nv/ComputeCache (size-capped, env-tunable).

The incident pattern every GPU platform hits once: new GPU generation lands in the fleet → no sm_XX cubin in anything → every process JITs every kernel on startup → 2–10 minute cold starts, CPU pegged, sometimes cache thrash if CUDA_CACHE_MAXSIZE is small. Fixes: add the new arch to build matrices (cost: binary size, build time), pre-warm the JIT cache in images, or lazy-load with CUDA_MODULE_LOADING=LAZY (CUDA 11.7+, now default — loads kernels on first use instead of module load).

As the platform owner you set this policy: which architectures get AOT SASS, which ride PTX JIT — a real cost/risk tradeoff (build farm minutes and gigabyte containers vs cold-start risk). Phase 10's release engineering inherits this decision; Phase 09's HAL generalizes it across vendors (AMD has no stable PTX equivalent — code objects are per-ISA, which changes the portability calculus entirely).

Chapter 8: What the Optimizer Does to Your Kernel

The transforms worth recognizing in Lab 01's dumps (and implementing, in miniature, in Lab 02's extensions):

  • FMA contraction: a*b + c → fma.rn — fewer instructions, different rounding (one rounding instead of two). -fmad=false disables it when bit-exactness vs CPU matters (you'll be asked this in numerical-debug contexts).
  • If-conversion / predication: short conditionals become predicated instructions — no branch, no divergence machinery; Phase 01 Ch. 4's "compiler often predicates" claim, now visible in SASS.
  • Loop unrolling: #pragma unroll or heuristics; buys ILP (Phase 01 Ch. 5's other latency-hiding axis) at the cost of registers and I-cache.
  • Constant folding / CSE / DCE: the bread-and-butter IR passes — your Lab 02 extension implements folding and CSE and measures emitted-instruction reduction.
  • Register allocation vs occupancy: ptxas balances registers-per-thread against spills; you nudge it with -maxrregcount/__launch_bounds__ and you measure rather than assume (Phase 02 Ch. 8's judgment call — this is where it's decided).

Misconception: "the compiler will fix my memory layout." It won't — coalescing (Phase 02 Ch. 6) is a data layout property the compiler cannot see past your pointer arithmetic. Compilers optimize instructions, not your AoS schema.

Chapter 9: The Portability Landscape — NVPTX, ROCm, SPIR-V, Triton, MLIR

The map of "hardware-independent GPU software" attempts — i.e., your competitive landscape:

  • LLVM NVPTX: upstream LLVM's PTX backend (what clang --cuda uses; NVIDIA's cicc is a private fork ahead of it). Matters because every "new language for GPUs" (Julia, Mojo, Rust-GPU, Triton) rides LLVM to PTX.
  • AMD ROCm/HIP: HIP is CUDA-API-shaped C++ (hipMalloc, hipLaunchKernel) — porting is mostly mechanical (hipify). The backend compiles to GCN/CDNA ISA directly — no stable virtual ISA layer: code objects target specific gfx90a/ gfx942 ISAs. Public ISA docs (unlike SASS!) but no PTX-like forward-compat contract — the key asymmetry in any NVIDIA/AMD platform strategy.
  • SPIR-V: Khronos' portable IR (Vulkan compute, OpenCL, SYCL). The truly vendor-neutral option; the performance ceiling for ML workloads has historically lagged, which is the whole tension: neutrality vs peak.
  • Triton: a Python DSL where you program at block granularity (tl.load a tile, compute, tl.store) and the compiler handles smem staging, vectorization, and auto-tuning — i.e., it automates Phase 02 Ch. 7. Lowered through MLIR/LLVM to PTX (and increasingly to AMD). Why it won developer mindshare: it sits at the right abstraction level — high enough to be productive, low enough to hit ~90% of hand-tuned performance on the kernels LLMs need.
  • MLIR: LLVM's "compiler for building compilers" — dialects (linalg, gpu, nvgpu) that progressively lower. The bet most ML-compiler teams (and most accelerator startups) are making for their middle layer.

The platform-strategy takeaway: there are exactly three credible positions for a hardware-agnostic platform: (a) target each vendor's kernel libraries and ship a HAL (fastest to market — Phase 09 builds this), (b) adopt Triton/MLIR as your kernel layer and inherit its backends, (c) own a full compiler (moonshot). Most real platforms blend (a) and (b). Be able to argue all three with costs.

Chapter 10: Compiler-as-Moat — The Business Reading

Why this phase belongs in a Head of Engineering track, in four observations:

  1. CUDA's dominance is a software annuity. PTX forward-compat (Ch. 4) means every kernel ever written keeps working — so ecosystem value compounds. Competitors must beat not just the chip but 18 years of accumulated, still-running software.
  2. Accelerator startups die in the middle layer. Silicon tape-out is hard but bounded; "PyTorch runs well" is unbounded — thousands of ops × precisions × shapes, each needing near-roofline kernels. Audit question for any vendor pitch: "show me your top-20 op coverage at p50/p99 of peak, on my shapes."
  3. Kernel velocity is a hiring/architecture choice you own. FlashAttention shipped as CUDA; within months every serious shop needed it across hardware. Platforms with a Triton/MLIR posture absorbed it in weeks; pure-CUDA shops on non-NVIDIA hardware waited quarters. Your Phase 09 HAL design decides which side of that your company lives on.
  4. The JD's "decouple software from hardware" is only as real as your compiler/kernel story. Anyone can wrap two SDKs behind an interface; the hard promise is performance-portable decoupling, and Chapters 4–9 are the technical content of that promise.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02. Read real PTX before generating your own — you'll plagiarize the prologue idioms, which is exactly the point.

Lab 01 (dissection):

  1. Annotate saxpy.ptx line-by-line using Ch. 5 as the rubric, before opening the answer key.
  2. SASS pass: find one predicated instruction, one FFMA, the unrolled loop body. With -maxrregcount 16: find LDL/STL spills.
  3. Fatbin dissection: cuobjdump --list-elf on a two-arch build; relate every entry to Ch. 7's load-time decision tree.
  4. No CUDA toolkit: godbolt.org → CUDA C++ → it shows PTX and SASS side by side; everything except the fatbin step works there.

Lab 02 (mini compiler):

  1. Run python solution.py first — see the five kernels compile and the differential tests pass; read one emitted PTX fully.
  2. Read in pipeline order (Ch. 2): Lexer → Parser → TypeChecker → CodeGen → PTXInterp. At each stage, ask "what's the input shape, what's the output shape" — they're all small list/tree transforms.
  3. Extensions in order: intrinsics (touch every stage once) → constant folding (your first IR pass) → CSE (your second; measure both) → the fused kernel.
  4. If you have any NVIDIA GPU: the README's cuda-python snippet loads your PTX for real. Do it — "my compiler's output runs on hardware" is a sentence few engineering leaders can say.

Success Criteria

  • You can draw the nvcc pipeline (5 boxes) and say what each tool does (Ch. 3)
  • You can read the saxpy PTX cold and narrate all five rubric points (Ch. 5)
  • You can spot spills, FMAs, and predication in a SASS dump (Ch. 6)
  • You can explain the fatbin decision tree and the cold-start incident + three fixes (Ch. 7)
  • Your compiler passes its differential tests and you've written at least one optimization pass (Ch. 2, 8)
  • You can argue the three platform-portability postures with costs (Ch. 9–10)

Interview Q&A

Q1: What happens, end to end, when you compile and run a CUDA program on a GPU that didn't exist when you compiled it? A: nvcc split the source; device code went C++ → NVVM/LLVM IR → PTX, and ptxas compiled PTX → SASS for each -code=sm_XX requested; fatbinary bundled the cubins plus (if requested) the PTX. On the new GPU there's no matching cubin, so the driver JIT-compiles the embedded PTX to the new architecture's SASS at module load, caching in the compute cache. First load is slow (the cold-start incident); subsequent loads hit the cache. If the fatbin contains no PTX (-code=sm_80 only), the load fails with "no kernel image available" — which is why production build policy always embeds a recent compute_XX.

Q2: Why does PTX exist at all — why not compile straight to machine code like CPUs do? A: CPU ISAs are stable contracts (x86 from 1978 still runs); GPU hardware ISAs change every generation because NVIDIA reserves the right to redesign encodings, scheduling, and register files for performance. PTX restores the missing stability: a virtual ISA with infinite typed registers and explicit state spaces that the driver can lower to any real architecture, past or future. It splits the compiler where you want it split — slow, smart optimization AOT into PTX; fast, mechanical lowering at load time. Strategically it's the CUDA ecosystem's compatibility promise — the moat (and AMD's lack of an equivalent is the key asymmetry in portability planning).

Q3: Reading a SASS dump, how do you tell register pressure became a problem? A: LDL/STL instructions — local loads/stores — are spills: values that didn't fit in registers round-tripping through local memory (which is HBM-backed, so each spill is a global-memory-latency event inside your inner loop). Cross-check with -Xptxas=-v ("N bytes spill stores/loads") and ncu's local-memory traffic. The causes: too many live values (deep unrolling, big per-thread arrays), or someone capped -maxrregcount/__launch_bounds__ chasing occupancy — the Phase 02 tradeoff where you measure both sides before deciding.

Q4: Why has Triton succeeded where earlier GPU DSLs struggled? A: Abstraction level. CUDA C++ is thread-level — maximum control, maximum effort; every tiling/staging decision is manual (Phase 02 Lab 01). Polyhedral/graph compilers (early TVM et al.) were too far above — schedules divorced from the memory hierarchy lose peak. Triton programs the block: you write tile loads and tile math; the compiler assigns threads, stages shared memory, vectorizes, and auto-tunes — automating exactly the mechanical parts of Phase 02 Ch. 7 while leaving algorithmic structure to the author. It hits ~90% of hand-tuned on the kernels that matter to LLMs, it's Python-embedded (the audience's language), and it rides MLIR/LLVM so backends beyond NVIDIA are tractable — which is also why it's a credible kernel layer for a hardware-agnostic platform.

Q5: Hand-compile c[i] = a[i] * 2.0 + b[i] to PTX. Sketch is fine. A: Prologue: .visible .entry with four .params (a, b, c pointers + n); ld.param each; cvta.to.global the pointers. Index: mov from %ctaid.x, %ntid.x, %tid.x; mad.lo.s32 %r4 = ctaid*ntid + tid. Guard: setp.ge.s32 %p1, %r4, n then @%p1 bra DONE. Body: mul.wide.s32 index → byte offset (×4), add.s64 to each base, ld.global.f32 %f1, [a_i], ld.global.f32 %f2, [b_i], then fma.rn.f32 %f3, %f1, 2.0, %f2, and st.global.f32 [c_i], %f3. DONE: ret. (Lab 02's compiler emits exactly this — and the interview answer is stronger for being able to say so.)

Q6 (leadership): A candidate for your compiler-team lead says "we should build our own IR." How do you evaluate? A: Probe for the build-vs-adopt analysis. MLIR exists precisely so teams don't build bespoke IRs: dialects, pass infrastructure, and existing lowerings to NVPTX/ROCm are free; a custom IR re-pays all of that for marginal fit gains. Legitimate custom-IR cases: a truly novel execution model the dialects can't express, or extreme compile-time constraints. I'd ask: which MLIR dialects did you evaluate and where exactly did they fail? What's the five-year maintenance headcount? How do you absorb the next FlashAttention-shaped kernel innovation (ecosystem velocity, Ch. 10)? A candidate with a crisp "we tried linalg+nvgpu, here's the specific gap" earns the conversation; "LLVM is bloated" doesn't.

References

  • NVIDIA, PTX ISA Specification — https://docs.nvidia.com/cuda/parallel-thread-execution/
  • NVIDIA, CUDA Compiler Driver NVCC (the pipeline doc; run nvcc -dryrun) — https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/
  • NVIDIA, CUDA Binary Utilities (cuobjdump, nvdisasm) — https://docs.nvidia.com/cuda/cuda-binary-utilities/
  • Compiler Explorer (Godbolt) CUDA mode — https://godbolt.org
  • Tillet, Kung, Cox, "Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations" (MAPL 2019); docs — https://triton-lang.org
  • Lattner et al., "MLIR: Scaling Compiler Infrastructure for Domain Specific Computation" (CGO 2021)
  • LLVM NVPTX backend guide — https://llvm.org/docs/NVPTXUsage.html
  • AMD ROCm/HIP documentation + public ISA guides (CDNA3) — https://rocm.docs.amd.com
  • Crafting Interpreters (Nystrom) — the gentlest serious compiler-construction text, for Lab 02 depth — https://craftinginterpreters.com
  • Cross-track: Phase 02 WARMUP Ch. 8 (occupancy/registers), Phase 09 WARMUP (the HAL this knowledge feeds)