« All Roles

Head of Software Engineering [GPU] — Complete Apprenticeship Curriculum

Target Roles:

• Head of Engineering / VP Engineering — AI Infrastructure (deep-tech startups)
• Principal / Staff Engineer — GPU Platforms, AI Compute Clouds (CoreWeave, Lambda, Together, Modal, NVIDIA, AMD)
• Platform Lead — LLM Serving & GPU Orchestration (vLLM-style platforms, inference clouds)
• Founding Engineer — Sovereign AI / On-Premise LLM Infrastructure
• Engineering Lead — Hardware-Agnostic AI Runtime / GPU Compiler Stack

Duration: 26–36 weeks core — extendable to 18 months for full head-of-engineering readiness Goal: Build the rare combination this JD demands: silicon-level GPU understanding, compiler and driver literacy, production LLM-serving depth (GPU scheduling, KV-cache management), hardware-decoupled platform architecture, security/licensing primitives for commercial software, and the leadership judgment to ship it all through a real release cycle.

Why This Curriculum Exists

The "Head of Software Engineering [GPU]" role is one of the hardest hiring profiles in the industry, because it sits at the intersection of five disciplines that are usually separate careers:

1. GPU systems engineering — chips, SMs, warps, memory hierarchies, CUDA, drivers, compilers down to PTX/SASS.
2. AI infrastructure — LLM serving, KV-cache management, continuous batching, multi-engine deployment (vLLM, TensorRT-LLM, SGLang).
3. Distributed systems & HPC — collectives, NCCL, multi-node scheduling, failure domains.
4. Platform productization — decoupling software from hardware (the JD's core mandate: "a platform that operates independently of the underlying hardware"), developer ecosystems, OEM/cloud integrations.
5. Commercial software protection — hardware fingerprinting, license enforcement, code obfuscation, air-gapped/sovereign deployment for regulated industries.

Most engineers know one or two of these. This curriculum is built backward from the JD so that every "required" and "nice to have" line maps to phases and hands-on labs:

What You Will Build

By the end of this curriculum you will have produced, with your own hands:

• A SIMT execution simulator that demonstrates warp divergence and lockstep execution (Phase 01)
• Memory-hierarchy microbenchmarks in C that measure cache lines, bandwidth ceilings, and the roofline of your own machine (Phase 01)
• CUDA kernels from vector-add to tiled shared-memory matmul, with occupancy/coalescing experiments (Phase 02)
• A mini compiler that emits real PTX from an expression language — plus annotated dissections of NVCC's pipeline, PTX, and SASS (Phase 03)
• A Linux character-device driver in C with an ioctl interface that mimics how GPU drivers expose command submission (Phase 04)
• A container-from-scratch lab (namespaces + cgroups) showing exactly how GPUs pass into containers (Phase 04)
• A caching GPU-style memory allocator in Rust (the PyTorch-allocator design) and a stream scheduler in C++ (Phase 05)
• A GPU cluster scheduler simulator (bin-packing, gang scheduling, fragmentation metrics) and a Kubernetes device-plugin walkthrough (Phase 06)
• A paged KV-cache allocator with prefix sharing and a continuous-batching simulator with TTFT/TPOT/goodput metrics (Phase 07)
• Ring all-reduce implemented over real sockets, benchmarked against the theoretical 2(N-1)/N bound (Phase 08)
• A hardware abstraction layer in C with dlopen-loaded vendor backends — the literal architecture of "decouple software from hardware" (Phase 09)
• A GPU metrics exporter (Prometheus format), failure-drill runbooks, and an SLO definition pack (Phase 10)
• A hardware fingerprinting tool, an Ed25519-signed license-enforcement system with grace periods and air-gap activation, and an anti-tamper/obfuscation lab in C (Phase 11)
• A capstone: mini GPU orchestration platform — scheduler + HAL + serving + licensing + observability composed into one system, with a design doc and release-cycle artifacts (Phase 12)

Folder Structure

Head of Software Engineering [GPU]/
├── README.md                                   ← You are here
├── jd.md                                       ← Target role profile
├── phase-01-gpu-architecture-silicon/          ← From transistor to SM: how the chip actually works
├── phase-02-cuda-programming/                  ← CUDA C++: kernels, memory, streams, occupancy
├── phase-03-gpu-compilers-codegen/             ← NVCC, PTX, SASS, LLVM NVPTX, Triton; build a mini PTX compiler
├── phase-04-linux-drivers-containers/          ← Kernel modules, ioctl, the GPU driver stack, cgroups, container GPU passthrough
├── phase-05-runtime-systems-memory/            ← Caching allocators, streams/events, async runtimes (Rust + C++)
├── phase-06-gpu-scheduling-virtualization/     ← MPS, MIG, time-slicing, k8s device plugins, cluster scheduling
├── phase-07-llm-serving-kv-cache/              ← KV-cache management, PagedAttention, continuous batching, multi-engine ops
├── phase-08-distributed-hpc-collectives/       ← NCCL, ring all-reduce, topology, multi-node failure domains
├── phase-09-hardware-abstraction-portability/  ← HAL design, plugin ABIs, dlopen backends — the decoupling mandate
├── phase-10-observability-production/          ← DCGM, exporters, eBPF, SLOs, release engineering
├── phase-11-security-licensing-protection/     ← Fingerprinting, license enforcement, obfuscation, attestation, air-gap
├── phase-12-leadership-capstone/               ← Head-of-engineering practice + mini orchestration platform capstone
├── interview-prep/                             ← Concepts, coding, systems, leadership interviews
└── system-design/                              ← GPU cloud, serving platform, HAL, licensing-at-scale walkthroughs

26-Week Core Schedule

Phase Structure

Every phase follows a consistent format:

Labs deliberately span the JD's language matrix: C (drivers, HAL, anti-tamper), C++/CUDA (kernels, stream scheduler), Rust (allocator, fingerprinting), Go (device plugin), Python (simulators, serving, licensing), TypeScript (platform CLI tooling in the capstone).

No GPU? You can still do ~85% of this track

Labs are designed in three tiers:

1. CPU-only (most labs): simulators, allocators, schedulers, compilers-to-PTX (text output), licensing, HAL — need only a Linux/macOS box.
2. Any-NVIDIA-GPU (Phase 02, parts of 03/10): a free Google Colab T4 or any consumer card works.
3. Cluster-scale (parts of 06/08): simulated locally; the WARMUPs explain what changes at real scale.

Prerequisites

• Solid Python; working C or C++ (you can read pointers, structs, makefiles)
• Comfort in a Linux shell (you'll compile kernels modules and use perf, strace)
• Basic ML literacy: what a transformer forward pass is (or run the LLM Inference Engineer track Phases 01–04 first)
• Git, Docker basics
• No prior CUDA, driver, or compiler experience required — Phase 01 starts at the silicon

Related Tracks in This Repo

• LLM Inference Engineer — deeper on models/training; its Phase 09 pairs with our Phase 07
• Security Engineer — deeper on offensive/defensive security; pairs with our Phase 11
• AI Specialist — air-gapped serving labs that pair with our Phase 11 sovereign-AI work

We’re partnering with a fast-scaling deep tech startup building next-generation AI infrastructure that combines cutting-edge hardware and software to deliver high-performance AI compute solutions.

The Opportunity

We’re hiring a Head of Engineering to lead the development of a platform that operates independently of the underlying hardware. This is a critical leadership role focused on building a developer-first system that enables flexibility, performance, and global adoption.

The Role

Lead the development of a standalone orchestration platform Define and execute the strategy to decouple software from hardware environments Build and scale a developer-focused platform and ecosystem Drive integrations with OEMs, cloud providers, and infrastructure partners Own platform adoption and commercialisation through partnerships Stay hands-on where needed—maintaining technical credibility with engineering teams Engage with customers and partners at both technical and executive levels

What They’re Looking For

Minimum 10+ years in software engineering with at least 4 years leading teams of 5 or more engineers. Deep expertise in at least two of: AI infrastructure, MLOps, distributed systems, GPU-accelerated workloads, HPC, or runtime systems programming Strong background in GPU scheduling, KV-cache management, or LLM serving optimisation Hands-on experience deploying and optimising AI infrastructure at production scale, ideally with multiple inference engines, utilities, and training frameworks. Proven ability to ship production software through a full release cycle, not just a prototype Strong fluency in Python, C++, and a systems language (Rust, Go, or modern C); comfort with JavaScript or TypeScript for tooling Solid grounding in Linux internals, containers, GPU drivers, CUDA, and observability tooling

Nice to Have

Direct experience with sovereign AI, on-premise LLM deployment, air-gapped systems, or AI in regulated industries (defence, energy, government, finance, healthcare) Open-source contributions in the AI infrastructure space Prior experience as a founding engineer or senior engineer at a deep-tech startup that reached commercial GA Experience with security primitives: hardware fingerprinting, license enforcement, software protection, code obfuscation

Phase 01 — GPU Architecture: From Silicon to SM

Difficulty: ⭐⭐⭐☆☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], GPU Platform Engineer, Performance Engineer Hardware needed: none — both labs run on any CPU

Why This Phase Exists

Every decision you will make as a GPU platform leader — scheduler design, KV-cache budgets, which inference engine to standardize on, what to tell an OEM partner their chip needs — reduces to a handful of physical facts about how GPUs are built. Leaders who can't derive "decode is memory-bandwidth-bound" from silicon get out-negotiated by vendors and out-designed by competitors.

This phase builds that base. Not trivia — mechanism. You will be able to explain, with numbers, why a GPU dedicates its transistor budget to ALUs instead of caches and branch predictors, why 32 threads execute in lockstep, why divergence halves throughput, and why "how many GB/s" matters more than "how many TFLOPs" for the workloads this JD cares about.

Concepts

• Why GPUs exist: latency-oriented vs throughput-oriented design; the transistor-budget argument
• The compute hierarchy: GPU → GPC → SM → warp scheduler → CUDA core / Tensor Core
• SIMT execution: warps, lockstep, active masks, divergence and reconvergence
• Latency hiding: zero-cost warp switching vs CPU context switches; occupancy as a latency-hiding budget
• The memory hierarchy: registers → shared memory/L1 → L2 → HBM; sizes, latencies, bandwidths for A100/H100
• HBM: why stacked DRAM, why bandwidth is the headline number
• Arithmetic intensity and the roofline model — the single most useful performance abstraction
• Tensor Cores: what an MMA unit actually computes; why FP16/BF16/FP8 exist
• NVLink, PCIe, and why interconnect shows up in Phase 08
• Reading a spec sheet critically: A100 vs H100 vs MI300X vs a startup accelerator

Labs

Lab 01 — SIMT Simulator (Python)

Lab 02 — Memory Hierarchy & Roofline Microbenchmarks (C)

Deliverables Checklist

• SIMT simulator runs; you can explain every line of the execution loop
• Divergence-cost table produced and explained
• membench results with identified cache plateaus
• Roofline worksheet: ridge points computed for your CPU and for A100/H100 from spec sheets
• You can do the "512 KB per token" style bandwidth math cold (see WARMUP Ch. 7)

Interview Relevance

• "Walk me through what happens when a warp hits a branch."
• "Why do GPUs have thousands of cores but no branch predictors?"
• "A100 has 1,555 GB/s HBM and 312 TFLOPS BF16 — what's the ridge point, and what does it mean for LLM decode?"
• "Your vendor claims their chip beats H100. What three numbers do you ask for?"

Warmup Guide — GPU Architecture: From Silicon to SM

Zero-to-expert primer for Phase 01. You need no prior hardware background: this guide starts at "what is a transistor budget" and ends with you computing rooflines for H100 and arguing with spec sheets. Every later phase stands on this one.

Table of Contents

• Chapter 1: Why GPUs Exist — The Transistor-Budget Argument
• Chapter 2: The Compute Hierarchy — GPU, GPC, SM, Warp
• Chapter 3: SIMT Execution — Warps, Lockstep, and the Active Mask
• Chapter 4: Divergence and Reconvergence
• Chapter 5: Latency Hiding — The GPU's Real Superpower
• Chapter 6: The Memory Hierarchy — Registers to HBM
• Chapter 7: Arithmetic Intensity and the Roofline Model
• Chapter 8: Tensor Cores and Reduced Precision
• Chapter 9: Interconnect — PCIe, NVLink, and Why It Matters Later
• Chapter 10: Reading a Spec Sheet Like a Head of Engineering
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: Why GPUs Exist — The Transistor-Budget Argument

Zero background: a chip is a fixed area of silicon. Every feature — arithmetic units, caches, control logic — costs transistors, and the budget is finite (H100: ~80 billion transistors on TSMC 4N). Architecture is the art of deciding what to spend them on.

The CPU's spend: a modern CPU core is mostly not arithmetic. It spends transistors on branch predictors, out-of-order execution windows, register renaming, prefetchers, and megabytes of cache. Why? Because a CPU's contract is minimize the latency of one instruction stream. All that machinery exists to keep a single thread from ever waiting.

The GPU's spend: a GPU inverts the contract: maximize the throughput of tens of thousands of threads, and let any individual thread wait as long as it likes. Under that contract, branch predictors and out-of-order windows are wasted transistors. Strip them out, and the same silicon area holds hundreds of ALUs instead of one fat core. The H100 has 16,896 FP32 lanes; a desktop CPU has on the order of hundreds (cores × SIMD width).

The catch: this trade is only a win if (a) the workload has massive data parallelism (the same operation over many elements — exactly what matmuls, convolutions, and attention are), and (b) you have a trick to stop stripped-down cores from stalling on memory. That trick is latency hiding (Chapter 5), and it is the deepest idea in this phase.

Production significance: this is why "port it to GPU" is not free. Code with unpredictable branches and pointer-chasing (graph traversal, parsers) hits the contract the GPU explicitly refused to optimize. When a partner asks "can your platform accelerate X," your first mental step is this chapter.

Common misconception: "GPUs are fast because they have more cores." No — per-lane, a GPU is slower than a CPU (lower clocks, in-order, no speculation). GPUs win only when parallelism and latency hiding apply. A single-threaded task runs ~10× slower on a GPU lane than on a CPU core.

Chapter 2: The Compute Hierarchy — GPU, GPC, SM, Warp

The physical organization, top-down, using H100 (SXM) numbers:

• GPU die → contains 8 GPCs (Graphics Processing Clusters) — mostly a layout/routing unit; you rarely reason about GPCs.
• GPC → contains SMs (Streaming Multiprocessors). H100 SXM: 132 SMs. The SM is the unit you reason about for almost everything: it is the GPU's "core."
• SM → contains 4 warp schedulers (sub-partitions), each owning:
  • 32 FP32 lanes ("CUDA cores") → 128 per SM
  • a slice of the register file (64K × 32-bit registers per SM — 256 KB; more than the L1!)
  • 1 Tensor Core per scheduler (4 per SM)
  • shared access to 228 KB of combined shared memory / L1
• Warp → 32 threads that execute as one unit (Chapter 3). Up to 64 warps resident per SM (2,048 threads).

Why this shape: the SM is a replication unit. NVIDIA scales product tiers by fusing off SMs (a defective die becomes a cheaper SKU) — which is why your platform code must never assume an SM count; it must query it (Phase 02's cudaGetDeviceProperties, Phase 09's capability negotiation).

The software mapping you'll meet in Phase 02: a CUDA thread block is assigned to exactly one SM and stays there; the grid of blocks spreads across SMs. Blocks are the unit of scheduling-onto-SMs; warps are the unit of execution-inside-SMs.

Misconception: "CUDA core" ≈ "CPU core." A CUDA core is a single FP32 ALU lane — no fetch, no decode, no program counter of its own. The closest CPU analogy to an SM is one CPU core with very wide SIMD and 64-way hyperthreading.

Chapter 3: SIMT Execution — Warps, Lockstep, and the Active Mask

SIMT = Single Instruction, Multiple Threads. The hardware fetches and decodes one instruction and issues it to 32 threads at once (a warp). One program counter per warp, not per thread (pre-Volta; see the nuance below).

Each thread has its own registers, so the same instruction operates on different data — thread 0 adds element 0, thread 31 adds element 31. From the programmer's view, you write scalar per-thread code; from the hardware's view, it's a 32-wide vector instruction.

The active mask: a 32-bit mask deciding which lanes commit results this instruction. All 32 lanes always march together; the mask says which ones' writes count. This is the entire mechanism behind divergence (next chapter), and it is literally what your Lab 01 simulator implements.

SIMT vs SIMD (favorite interview distinction):

• SIMD (CPU AVX): you (or the compiler) must pack data into vector registers and handle edges/masks explicitly. Vector width is in the ISA.
• SIMT: you write scalar code per thread; the hardware handles masking and the warp width is a hardware property. Divergence is legal (just slow), pointers can differ arbitrarily per lane. SIMT = SIMD with per-lane control flow and addressing, managed by hardware.

Volta+ nuance — Independent Thread Scheduling: since Volta, each thread has its own PC architecturally, and the hardware interleaves divergent paths, which makes intra-warp synchronization deadlock-free (and is why __syncwarp() and the _sync suffixed intrinsics exist). Execution is still mask-serialized — divergence still costs the same — but correctness semantics changed. Saying "one PC per warp, and since Volta that's virtualized per-thread" marks you as current.

Chapter 4: Divergence and Reconvergence

What happens at if (threadIdx.x < 16) A(); else B(); inside one warp:

1. Lanes 0–15 want path A; lanes 16–31 want path B.
2. Hardware executes A with mask 0x0000FFFF, then B with mask 0xFFFF0000.
3. At the merge point the mask returns to 0xFFFFFFFF — reconvergence (classically tracked with a stack of (mask, reconvergence-PC) entries — your Lab 01 extension implements exactly this).

The cost model: time = sum of the divergent paths' lengths, not the max. Two equal halves → 2× slowdown. Worst case, all 32 lanes take different paths (a 32-way switch on threadIdx.x % 32) → 32× slowdown: each instruction issues 32 times with one live lane. Lab 01 makes you measure this, not memorize it.

What does NOT diverge (the production rule of thumb):

• Branches uniform across the warp (if (blockIdx.x % 2)) — every lane agrees, one path executes. Free.
• Branches on data that happens to agree per-warp — also free at runtime.
• Short divergent paths — the compiler often converts them to predication (execute both sides, mask the writes), trading a few wasted instructions for no branch machinery.

Production significance: this is why GPU code sorts/buckets work by type before processing (e.g., batching same-length sequences in serving — Phase 07 inherits this idea), and why % 32-pattern branches are a code-review red flag in kernels.

Misconception: "divergence breaks correctness." Never — masks preserve semantics. Divergence is purely a throughput tax.

Chapter 5: Latency Hiding — The GPU's Real Superpower

The deepest single idea in GPU architecture.

The problem: HBM access latency is ~400–800 cycles. A CPU hides this with caches + prefetching + out-of-order execution — all the transistor-expensive machinery GPUs deleted. So how does a GPU not stall constantly?

The answer — oversubscription: each SM holds up to 64 resident warps but only issues from ~4 per cycle. When warp 7 issues a load, the scheduler doesn't wait: next cycle it issues from warp 23, then warp 41... By the time the data arrives 400 cycles later, warp 7 simply rejoins the eligible pool.

Why switching is free: on a CPU, a context switch saves/restores registers to memory (microseconds). On an SM, all 64 warps' registers live in the register file simultaneously — that's why it's 256 KB, bigger than L1. A "switch" is the scheduler picking a different index. Zero cycles.

Occupancy = resident warps / 64. It's a latency-hiding budget: with 400-cycle stalls and ~4-cycle compute between them, you need enough other warps to fill 400 cycles of issue slots. What limits occupancy: registers per thread (64K regs ÷ 256 regs/thread = 256 threads = 8 warps — bad), shared memory per block, block size. You'll tune this for real in Phase 02 Lab 02.

Misconception: "100% occupancy = fastest." No — occupancy only needs to be high enough to cover latency. A kernel with high arithmetic intensity may run at peak at 25% occupancy; pushing occupancy by cutting registers-per-thread can spill registers to memory and slow it down. Volkov's "Better Performance at Lower Occupancy" is the canonical reference.

Chapter 6: The Memory Hierarchy — Registers to HBM

Numbers to know cold (H100 SXM; A100 in parens where it differs meaningfully):

Walk the gradient: each level down is ~10× bigger and ~3–10× slower/narrower. The two cliffs that dominate everything: HBM↔L2 (the "memory wall" inside the device) and device↔host over PCIe (50× narrower than HBM — why "keep data on the GPU" is rule #1, and why Phase 05's allocator exists to avoid round trips).

Shared memory deserves emphasis: it is software-managed L1 — a scratchpad the kernel explicitly loads. The tiled matmul (Phase 02 Lab 01) stages tiles in shared memory so each HBM byte is reused ~tile-size times. Bank conflicts (32 banks, 4-byte wide) are the shared-memory analog of divergence: conflicting lanes serialize.

HBM in one paragraph: DRAM dies physically stacked next to the GPU on a silicon interposer, connected by thousands of through-silicon vias → an extremely wide bus (5,120 bits vs 64 for DDR) at modest clocks. High bandwidth, still-high latency — bandwidth is what you can buy with parallelism; latency you hide (Chapter 5).

Misconception: "the GPU's 80 GB is like RAM, the model just fits or doesn't." Capacity is only the first constraint; bandwidth is the performance constraint (Chapter 7), and fragmentation of that capacity is a real engineering problem (Phase 05 allocator, Phase 07 paged KV).

Chapter 7: Arithmetic Intensity and the Roofline Model

The one-chart performance model you will use weekly for the rest of this track.

Arithmetic intensity (AI) = FLOPs performed ÷ bytes moved from memory. It is a property of the algorithm, not the hardware.

• vector_add: 1 FLOP per 12 bytes (read a, read b, write c, FP32) → AI ≈ 0.08
• naive matmul (no reuse): AI ≈ O(1)
• tiled/blocked matmul: AI grows with tile size — large GEMMs reach AI in the hundreds
• LLM decode at batch 1: ~2 FLOPs per weight byte → AI ≈ 2 (with FP16 weights, ~1)

The roofline: attainable FLOP/s = min(peak compute, AI × memory bandwidth). Plot: x = AI (log), y = FLOP/s (log). Flat roof = compute peak; slanted roof = bandwidth × AI. The ridge point = peak ÷ bandwidth is where they meet:

$$AI_{ridge} = \frac{\text{peak FLOP/s}}{\text{memory bandwidth}}$$

H100 BF16 dense: 989e12 ÷ 3.35e12 ≈ 295 FLOP/byte. A100: 312e12 ÷ 2.0e12 ≈ 156.

The consequence that funds this entire JD: LLM decode sits at AI ≈ 2, two orders of magnitude left of the ridge → decode is bandwidth-bound, GPU compute sits ~99% idle at batch 1, and every serving technique in Phase 07 (batching, quantization, paging, speculation) is an attempt to move right on this chart or to buy more of the slanted roof. The 7B-model arithmetic: 14 GB of FP16 weights ÷ 3.35 TB/s ≈ 4.2 ms/token ≈ 240 tok/s ceiling at batch 1 — regardless of TFLOPs.

How to use it as a leader: before any optimization, compute the kernel's AI and place it. Memory-bound → quantize, fuse, batch, improve locality; compute-bound → Tensor Cores, lower precision, better tiling. Optimizing FLOPs on a bandwidth-bound kernel is a category error you will now never make, nor accept in a design review.

Misconception: comparing chips by TFLOPs. For inference fleets, bandwidth and capacity usually decide; the spec-sheet TFLOPs figure often assumes sparsity ("with 2:4 structured sparsity" — read footnotes) and a precision your workload may not use.

Chapter 8: Tensor Cores and Reduced Precision

What a Tensor Core is: a fixed-function unit that computes a small matrix multiply-accumulate, e.g. D = A×B + C on 16×16-ish fragments, in one instruction (mma.sync at the PTX level — you'll see it in Phase 03). One H100 Tensor Core does hundreds of FLOPs/cycle vs 2 for a fused multiply-add on a CUDA core. That's how the spec jumps from ~67 TFLOPS FP32 to ~989 TFLOPS BF16.

Why precisions multiply (FP16/BF16/FP8/INT8): halving the element size doubles both effective bandwidth (Chapter 7's slanted roof) and Tensor Core throughput.

• FP16: 1 sign + 5 exp + 10 mantissa — precise but narrow range (overflows at 65,504; why training needed loss scaling)
• BF16: 1 + 8 + 7 — FP32's range, less precision; the training default
• FP8 (E4M3/E5M2, Hopper+): inference and now training, with per-tensor scaling
• INT8/INT4: quantized inference (Phase 07 / llm-inference track Phase 09)

The fine print that matters in vendor talks: Tensor Cores only engage when data layout, precision, and dimensions match their fragment shapes — which is why cuBLAS/cuDNN pad dimensions to multiples of 8/16, and why "your model got faster by padding the vocab size to a multiple of 64" is a real production anecdote.

Chapter 9: Interconnect — PCIe, NVLink, and Why It Matters Later

A single H100 reads its own HBM at 3.35 TB/s, talks to peers over NVLink at 900 GB/s (bidirectional, 18 links), and to the host over PCIe 5.0 at ~64 GB/s. Memorize the ratio: HBM : NVLink : PCIe ≈ 50 : 13 : 1.

Consequences you'll meet again:

• Tensor parallelism (splitting one model across GPUs) is only viable over NVLink-class links — the per-layer all-reduces die over PCIe (Phase 08).
• "Just offload the KV cache to host RAM" costs 50× bandwidth — viable only for cold data (Phase 07's swapping tier).
• Multi-node = InfiniBand/RoCE at ~400 Gb/s ≈ 50 GB/s — another order down; topology-aware scheduling (Phase 06) exists because of this gradient.

Chapter 10: Reading a Spec Sheet Like a Head of Engineering

The checklist when a vendor (or an OEM partner — this JD's daily life) hands you a datasheet:

1. Memory bandwidth and capacity first (inference fleets live and die here).
2. Sustained vs peak: is the TFLOPs number "with sparsity"? At what clocks (boost vs sustained under power cap)? Dense BF16 is the honest comparable.
3. Ridge point: compute it; compare against your workload's AI.
4. Interconnect: peer-to-peer bandwidth and topology, not just "has NVLink."
5. Software: what's the kernel library story? A chip without a mature compiler
   • kernels (Phase 03) delivers a fraction of paper FLOPs. This is the question for startup accelerators — and the gap your hardware-agnostic platform (Phase 09) monetizes.
6. Power/TCO: tokens per joule increasingly decides procurement.
7. MIG/virtualization support (Phase 06): can you sell fractions of it?

Run this checklist on A100 vs H100 vs MI300X as the Chapter exercise — MI300X's larger HBM (192 GB) vs H100's software maturity is the canonical real-world tradeoff discussion.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02. Lab 01 cements Chapters 3–5 (execution); Lab 02 cements Chapters 6–7 (memory). Don't invert: divergence is easier to internalize first.

Lab 01 (SIMT simulator):

1. Read solution.py top-down before running; predict what the divergence report will show.
2. Run; check predictions. The built-in asserts verify lockstep semantics.
3. Extensions in order: (a) reconvergence stack — the (mask, pc) discipline from Chapter 4; (b) vote.any — your first warp-cooperative primitive; (c) write the maximally divergent kernel and predict 32× before measuring.

Lab 02 (membench):

1. make && ./membench — then stare at the latency curve until you can point at L1, L2, L3, DRAM.
2. Fill in ROOFLINE.md: your CPU's ridge point from measured bandwidth and estimated peak; then A100's and H100's from spec sheets.
3. The capstone question: place batch-1 LLM decode (AI≈2) on the H100 roofline and write three sentences on what that implies — those three sentences are Phase 07's thesis.

Success Criteria

You're done with this phase when — without notes:

• You can explain the transistor-budget argument in 60 seconds (Ch. 1)
• You can define warp, active mask, and reconvergence, and state the 2×/32× divergence cost cases (Ch. 3–4) — and your simulator's output proves it
• You can explain why warp switching costs zero cycles and what occupancy actually buys (Ch. 5)
• You can reproduce the memory-hierarchy table within 2× from memory and identify your own machine's plateaus in your measured curve (Ch. 6)
• You can compute a ridge point and the 7B-decode token ceiling on any GPU given two spec numbers (Ch. 7)
• You can name the three spec-sheet questions that expose a weak accelerator pitch (Ch. 10)

Interview Q&A

Q1: Why don't GPUs have branch predictors? A: A branch predictor exists to keep a single instruction stream from stalling — latency optimization. The GPU's contract is throughput across thousands of threads: when a warp stalls (on memory or an unresolved branch), the scheduler issues from another resident warp at zero cost, because all warps' registers live permanently in the register file. Latency is hidden by parallelism rather than removed by prediction, so prediction hardware would be transistors taken away from ALUs for a problem the architecture already solves. Same argument deletes out-of-order execution and large caches.

Q2: A warp executes if (threadIdx.x % 2) A(); else B(); — what happens, exactly? A: One warp-level PC walks both paths serially. Hardware pushes the reconvergence point, executes A with active mask 0xAAAAAAAA, then B with 0x55555555, then reconverges to full mask. Both paths' instruction counts add → ~2× slowdown; results are still correct because masked lanes don't commit writes. If A and B are short, the compiler instead predicates (issues both sides, masks writes) and avoids branch overhead. Had the condition been uniform per warp (blockIdx.x % 2), there'd be no divergence at all. Post-Volta, threads have independent PCs architecturally (deadlock-free sync), but the serialization cost is unchanged.

Q3: Compute H100's ridge point and tell me what it means for LLM serving. A: Ridge = peak FLOP/s ÷ bandwidth = 989e12 (dense BF16) ÷ 3.35e12 B/s ≈ 295 FLOP/byte. Batch-1 decode performs ~2 FLOPs per weight byte read (AI ≈ 1–2), about 150× below the ridge → decode is purely bandwidth-bound; the ceiling for a 14 GB FP16 7B model is 3.35 TB/s ÷ 14 GB ≈ 240 tok/s with compute ~99% idle. Therefore serving economics come from raising effective AI: batching (reuse each weight byte across B requests), quantization (fewer bytes per weight), speculative decoding (more tokens per weight pass). That chain of reasoning is the foundation of every Phase 07 technique.

Q4: What's the difference between SIMT and SIMD, and why did NVIDIA choose SIMT? A: SIMD exposes vector width in the ISA: software packs registers, handles edge masks, can't have per-lane control flow or scattered addressing without explicit gather/scatter. SIMT presents a scalar program per thread; hardware groups 32 threads into a warp, runs them in lockstep, and manages masks automatically on divergence — per-lane branching and arbitrary per-lane addresses are legal, just potentially slow. NVIDIA chose SIMT because it makes the programming model scale: the same scalar kernel runs on any warp width or SM count, which is also why CUDA code from 2008 still runs on Hopper. The cost is a hardware tax (mask/reconvergence tracking) and the foot-gun that naive branchy code silently serializes.

Q5: When is increasing occupancy the wrong optimization? A: Occupancy is a latency-hiding budget, not a goal. If a kernel already has enough resident warps to cover its stall latency (or is compute-bound with high ILP), more occupancy adds nothing — and the usual way you get more occupancy (fewer registers per thread via -maxrregcount or __launch_bounds__) can force register spills to local memory, adding the very memory traffic you were hiding. Volkov's classic result: some kernels hit peak at ~25% occupancy via instruction-level parallelism. The discipline: profile first (Phase 02), identify whether stalls are actually latency-bound, and treat occupancy as one lever among register count, tiling, and ILP.

Q6 (leadership): An OEM partner's new accelerator claims "2× H100 performance." How do you evaluate the claim for your platform? A: Decompose the claim along the roofline: 2× on what — dense BF16 TFLOPs, HBM bandwidth, or an app benchmark with sparsity footnotes? For our serving-dominated workloads I'd ask for: dense (non-sparse) throughput per precision, memory bandwidth and capacity, peer-to-peer interconnect topology, power envelope, and — most decisive — the software story: compiler maturity, kernel library coverage (attention/GEMM), and graph-capture support, because paper FLOPs without kernels deliver a fraction of peak (the exact gap our hardware-abstraction layer, Phase 09, exists to bridge). Then I'd run our own AI-binned benchmark suite: a bandwidth-bound decode workload, a compute-bound prefill/GEMM workload, and a collective-heavy multi-chip workload, and compare tokens/joule and tokens/$ rather than TFLOPs.

References

• Hennessy & Patterson, Computer Architecture: A Quantitative Approach, 6th ed. — Ch. 4 (data-level parallelism, GPUs)
• NVIDIA, CUDA C++ Programming Guide — Hardware Implementation chapter — https://docs.nvidia.com/cuda/cuda-c-programming-guide/
• NVIDIA H100 Architecture Whitepaper — https://resources.nvidia.com/en-us-tensor-core
• NVIDIA A100 Architecture Whitepaper ("Ampere") — https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/nvidia-ampere-architecture-whitepaper.pdf
• Williams, Waterman, Patterson, "Roofline: An Insightful Visual Performance Model" (CACM 2009)
• Volkov, "Better Performance at Lower Occupancy" (GTC 2010) — https://www.nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf
• Jia et al., "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" (arXiv:1804.06826)
• Lindholm et al., "NVIDIA Tesla: A Unified Graphics and Computing Architecture" (IEEE Micro 2008) — the original SIMT paper
• AMD CDNA3 (MI300) Whitepaper — for the cross-vendor comparison exercise
• Cross-track: LLM Inference Engineer Phase 09 WARMUP — the serving-side consequences of Ch. 7

🛸 Hitchhiker's Guide — Phase 01: GPU Architecture

Read this if: you can write software but a GPU is still a black box that "does matmuls fast." This guide is the practitioner's compressed tour — the numbers, mental models, and war stories that the WARMUP derives from first principles.

0. The 30-second mental model

A GPU is a throughput machine: ~100+ small in-order cores (SMs), each running up to 64 groups of 32 lockstep threads (warps), switching between them every cycle for free to hide ~500-cycle memory latency. Performance = min(compute peak, arithmetic intensity × bandwidth) — the roofline. For the workloads this JD cares about (LLM decode), AI ≈ 2 and the H100 ridge is ≈ 295, so bandwidth is the product and TFLOPs are marketing.

If you remember one sentence: GPUs trade single-thread latency machinery for ALUs, and hide memory latency with warp oversubscription instead of caches.

1. The org chart of an H100

GPU (1 die, ~80B transistors)
└── 132 SMs                          ← "cores"; the unit you reason about
    └── 4 warp schedulers per SM     ← each issues 1 instr/cycle from eligible warps
        ├── 32 FP32 lanes            ← "CUDA cores" = ALU lanes, nothing more
        ├── 1 Tensor Core            ← small-matrix MMA unit; where the TFLOPs live
        └── register file slice      ← 256 KB/SM total; all resident warps live here
    └── 228 KB shared mem / L1       ← software-managed scratchpad
└── 50 MB L2 (device-wide)
└── 80 GB HBM3 @ 3.35 TB/s
└── 18 NVLink4 links @ 900 GB/s aggregate; PCIe 5 @ 64 GB/s

Bandwidth ratio to tattoo on your arm: HBM : NVLink : PCIe ≈ 50 : 13 : 1.

2. Execution: warps and divergence in 4 bullets

• 32 threads = 1 warp = 1 instruction stream. Active mask decides who commits.
• Divergent branch → paths run serially with complementary masks → cost is the sum of path lengths. 2-way even split = 2×; pathological 32-way = 32×.
• Uniform branches (same decision warp-wide) are free. blockIdx-based: free. threadIdx % 2: 2×. Code review accordingly.
• Post-Volta threads have independent PCs (sync inside divergent code won't deadlock), but the serialization cost is identical. Use __syncwarp() / *_sync intrinsics.

3. Latency hiding beats caching

The numbers that make the design click:

Occupancy = resident-warp fraction. It's a budget for hiding stalls, not a score: plenty of kernels peak at 25–50% occupancy (Volkov). The classic failure: chasing 100% occupancy with -maxrregcount, causing register spills, getting slower.

4. The memory wall, quantified

Per-byte energy and latency dominate compute by orders of magnitude. Practical table (H100):

Rules that fall out: keep data resident (Phase 05 allocator), stage reuse through shared memory (Phase 02 tiled matmul), batch host↔device transfers and overlap with compute (streams), never put a synchronous PCIe copy in a decode loop (a real production incident pattern — it shows up as mysterious 5× TPOT regressions).

5. Roofline: the only performance chart you need

attainable = min(peak_flops, AI × bandwidth),  AI = FLOPs / bytes

Workflow: compute AI → place on roofline → choose lever (memory-bound: quantize, fuse, batch, cache; compute-bound: precision, Tensor Cores, tiling). Refuse to review any "optimization" PR or vendor claim that skips this step.

6. Tensor Cores & precision menu

• Tensor Core = fixed-function fragment MMA (mma.sync in PTX). FP32 path: 67 TFLOPS. Tensor Core BF16: 989. The headline number is only reachable through them — via cuBLAS/cuDNN/CUTLASS or Triton, not hand-rolled FP32 loops.
• Precision ladder: FP32 → TF32 (free Ampere training win) → BF16 (training default; FP32 range) → FP16 (more mantissa, less range; loss-scaling era) → FP8 E4M3/E5M2 (Hopper inference/training) → INT8/INT4 (quantized inference). Each halving doubles bandwidth-limited throughput AND Tensor Core rate — a double roofline win, which is why quantization is the most leveraged inference optimization (Phase 07).
• Gotcha: dimensions must align to fragment shapes; padding to multiples of 8/16/64 is sometimes a real speedup. ("We made the vocab a multiple of 64 and gained 8%." — true story across several labs.)

7. Spec-sheet self-defense (the leadership skill)

Vendor says "2× H100." You ask:

1. Dense or sparse TFLOPs? (2:4 sparsity footnote inflates 2×.)
2. Sustained clocks under power cap, or boost?
3. HBM bandwidth and capacity? (Inference fleets are bought on these.)
4. P2P topology — all-to-all NVLink-class, or PCIe islands?
5. What runs on it today? Compiler maturity, attention/GEMM kernel coverage, framework integration. Silicon without software delivers ~20–40% of paper peak for the first years (every new accelerator's story; also the market gap a hardware-agnostic platform exploits).
6. Tokens/joule on your benchmark mix — AI-binned: one bandwidth-bound, one compute-bound, one collective-heavy.

MI300X vs H100 is the live case study: 192 GB vs 80 GB HBM (capacity win → fewer GPUs per model), comparable bandwidth, but a software-ecosystem gap that erodes the paper advantage. Where would your platform's HAL (Phase 09) change that calculus? That question is essentially this JD's pitch.

8. What you should be able to do after this phase

• Sketch the H100 org chart from memory with sizes/bandwidths within 2×.
• Derive the 240 tok/s 7B ceiling in 30 seconds from two spec numbers.
• Predict divergence cost of a code snippet by inspection.
• Run a roofline argument in a design review — and detect when someone else's performance claim violates one.
• Interrogate a spec sheet down to its footnotes.

Next: Phase 02 puts your hands on the machine — you'll write the kernels these mental models describe, and Nsight will show you warps stalling exactly where this guide says they will.

Lab 01 — SIMT Simulator: Warps, Masks, and Divergence

Phase: 01 — GPU Architecture | Difficulty: ⭐⭐⭐☆☆ | Time: 3–4 hours Language: Python (stdlib only) | Hardware: none

Concept primer: ../WARMUP.md Ch. 3–4 (SIMT, active mask, divergence), ../HITCHHIKERS-GUIDE.md §2.

Run

python solution.py        # runs the simulator + divergence benchmark + asserts

0. The mission

Build (and then extend) a tiny SIMT machine: one program counter per warp, 32 lanes, an active mask — and use it to measure what divergence costs instead of taking the textbook's word for it.

The two questions you must be able to answer at the end:

1. Why does if (tid % 2) cost 2× but if (tid < 32) cost nothing (per warp)?
2. Why is the worst case exactly 32× — and what real code patterns approach it?

1. The ISA

The simulator executes a list of instructions over 32 lanes. Each lane has its own register file; the warp has one PC and a mask stack:

This mirrors real hardware closely enough to be honest: PTX has @p bra (predicated branch) and the hardware tracks reconvergence via a stack of (mask, pc) entries (classic pre-Volta SSY/SYNC discipline — see WARMUP Ch. 3 for the Volta+ nuance).

2. How divergence is executed

When BRZ finds that some active lanes branch and some don't:

1. Push (current_mask, reconvergence_pc) onto the mask stack.
2. Run the fall-through path with the not-taken lanes' mask.
3. At RECONV, switch to the taken path with the complementary mask.
4. At RECONV again, pop the stack — full mask restored, paths rejoined.

Every instruction executed counts once per issue, regardless of how many lanes were active — that's the physical cost model: an instruction with 1 live lane costs the same issue slot as one with 32.

3. What the benchmark shows

solution.py runs the same workload four ways and reports issue counts:

Expected output (yours will match exactly — it's deterministic):

kernel            issues   util%  slowdown
uniform                6   100.0     1.00x
warp_uniform           7   100.0     1.17x
even_odd              10    70.0     1.67x
max_divergent        226    58.4    37.67x

util% = lane-cycles doing committed work ÷ (issues × 32) — the simulator's version of Nsight's "warp execution efficiency" metric you'll meet in Phase 02.

Two readings worth pausing on:

• even_odd shows 1.67×, not 2.0× — because only the divergent region doubles; the shared prologue (TID/MOD) issues once. Real kernels behave the same way: divergence cost is proportional to how much code sits inside the split. Isolate the divergent region and you'll measure the clean 2×.
• max_divergent utilization is 58%, not 3% — the 32-way dispatch chain (compare/branch per case) runs fully populated; only each one-lane body wastes 31 lanes. The slowdown is still catastrophic (37×). This mirrors real switch lowering on GPUs.

4. Extension exercises (do them — this is the lab)

1. Reconvergence stack depth: nest two divergent branches and assert the stack reaches depth 2 and unwinds correctly. (The skeleton handles one level; nested divergence is where naive implementations break.)
2. VOTE_ANY rd, ra: warp-cooperative primitive — rd = 1 in all active lanes if any active lane has ra != 0. This is __any_sync(). Then explain in a comment why warp votes need no memory traffic on real hardware.
3. Predication: add SELP rd, ra, rb, rp (lane-wise select) and rewrite even_odd without branches. Compare issue counts — you've just reproduced the compiler's if-conversion decision (WARMUP Ch. 4).
4. Latency hiding mini-model: simulate 4 resident warps where LD stalls a warp for 40 cycles; a scheduler issues from any non-stalled warp each cycle. Measure utilization with 1, 2, 4 warps and observe the occupancy curve flatten — Chapter 5 made quantitative.

5. Common pitfalls

1. Forgetting to count issues for instructions with few active lanes — the whole point is that issue cost is mask-independent.
2. Reconverging at the branch target instead of the post-dominator — your even_odd will produce wrong values in half the lanes.
3. Letting inactive lanes commit writes — masks must gate writes, not reads.
4. Treating JMP inside a divergent region as a reconvergence — it isn't; only RECONV pops.

6. What this lab proves about you

You can explain SIMT execution mechanically — masks, stacks, issue costs — which is the difference between "GPUs don't like branches" hand-waving and the engineering-leader version: "this kernel's branch is warp-uniform, it's free; that one is per-lane data-dependent and our profile shows 41% warp efficiency, here's the restructuring." Phase 02's profiler sessions will show you these exact numbers on real silicon.

Lab 02 — Memory Hierarchy & Roofline Microbenchmarks (C)

Phase: 01 — GPU Architecture | Difficulty: ⭐⭐⭐☆☆ | Time: 3–5 hours Language: C11 + make | Hardware: any CPU (the methodology transfers to GPU)

Concept primer: ../WARMUP.md Ch. 6–7 (memory hierarchy, roofline), ../HITCHHIKERS-GUIDE.md §4–5.

Run

make
./membench            # latency curve + bandwidth + roofline placement

0. The mission

Stop trusting tables — measure a memory hierarchy yourself, with the same two techniques every GPU microbenchmarking paper uses (e.g., Jia et al.'s Volta dissection):

1. Pointer-chasing measures latency: each load's address depends on the previous load's value, so the CPU cannot overlap or prefetch — you see the raw round-trip per level.
2. Streaming measures bandwidth: independent sequential accesses let the hardware pipeline everything — you see the peak transfer rate.

Then you place two kernels on your machine's roofline and verify the model predicts which one is memory-bound. This is the exact methodology you'd use to characterize an unknown accelerator a vendor hands you (WARMUP Ch. 10).

1. What the program does

• Latency sweep: builds a randomly-permuted cycle of pointers in buffers from 4 KB to 256 MB and chases it 10M times. Plotting ns/load vs buffer size shows plateaus at each cache level, with cliffs between them.
• Bandwidth: streams a large array (read-sum, then memcpy-style write) and reports GB/s.
• Roofline placement: times vector_add (AI ≈ 0.083 FLOP/byte) and a blocked matmul (AI grows with block size), converts to GFLOP/s, and prints where each lands relative to min(peak, AI × measured_bandwidth).

Example output (Apple M-series; your numbers will differ — that's the point):

== latency (pointer chase) ==
     4 KB     1.2 ns/load     <- L1
   128 KB     3.1 ns/load     <- L2
     8 MB     9.8 ns/load     <- SLC/L3
   256 MB    98.4 ns/load     <- DRAM
== bandwidth (stream) ==
read  : 61.2 GB/s   write : 38.9 GB/s
== roofline ==
vector_add: AI=0.083  predicted ceiling  5.1 GFLOP/s  measured  4.7  -> memory-bound (92% of roof)
matmul b=64: AI≈10.7  measured 41.3 GFLOP/s           -> compute-bound region

2. Fill in ROOFLINE.md

The worksheet asks you to:

1. Identify each plateau and match it to your CPU's published cache sizes.
2. Compute your CPU's ridge point from measured bandwidth + estimated peak (cores × clock × SIMD width × 2 FLOPs).
3. Compute A100 and H100 ridge points from their spec sheets (WARMUP Ch. 7 has the answers — derive before checking).
4. Place batch-1 LLM decode (AI ≈ 2) on the H100 roofline and write the three-sentence implication (this is Phase 07's thesis statement).

3. Why the methodology transfers to GPUs

4. Common pitfalls

1. Sequential pointer chase — the prefetcher hides DRAM latency and you'll "measure" L1 everywhere. The permutation must be random. (Try it: replace the shuffle with identity and watch the curve flatten — instructive!)
2. Compiler deleting your loop — the chase result must feed a volatile sink or the optimizer removes everything. Same trick JMH uses (Blackhole), same trick you'll need in CUDA benchmarks.
3. Timing one pass — too noisy; the harness does warmup + multiple reps and reports the minimum (least-disturbed) run.
4. Forgetting -O2 — at -O0 you're benchmarking the compiler's laziness, not the memory system.

5. Extensions

• Add a strided read benchmark (stride 1, 2, 4, … 128 elements) and find the cliff at your cache-line size (64 B): past it, every load is a new line and bandwidth collapses by the stride factor. This is exactly the GPU coalescing argument of Phase 02 — uncoalesced access wastes the same way.
• Run under taskset -c 0 vs a core on another socket/cluster (Linux) and observe NUMA / cluster effects.
• Use perf stat -e cache-misses,cache-references ./membench and correlate the counters with your plateaus.

6. What this lab proves about you

You can characterize any memory system — CPU today, an OEM partner's accelerator next quarter — with two from-scratch benchmarks and place workloads on its roofline before a single line of vendor marketing reaches your team.

Phase 02 — CUDA Programming: Kernels, Memory, Streams

Difficulty: ⭐⭐⭐⭐☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], GPU Platform Engineer, Kernel/Performance Engineer Hardware needed: any NVIDIA GPU — a free Google Colab T4 is enough for every lab here

Why This Phase Exists

The JD says "stay hands-on where needed — maintaining technical credibility with engineering teams." For a GPU platform, that credibility has one test: can you write, debug, and profile a CUDA kernel yourself? Not because the Head of Engineering writes kernels daily, but because every estimate you approve, every "it's a kernel problem" escalation you arbitrate, and every hire you evaluate depends on having done it.

Phase 01 gave you the machine model. This phase makes you operate the machine: the CUDA programming model (grids/blocks/threads), the memory API and its async variants, shared-memory tiling, coalescing, occupancy, streams, and the profiler that tells you the truth.

Concepts

• The CUDA software stack: runtime API vs driver API; what nvcc produces (handoff to Phase 03)
• Kernel launch anatomy: grid/block/thread, <<<blocks, threads, smem, stream>>>, dimension math
• The block-to-SM contract: why blocks must be independent; what that buys the hardware
• Error handling discipline: cudaGetLastError, sync vs async errors, the CUDA_CHECK macro
• Memory management: cudaMalloc/cudaMemcpy, pinned (page-locked) memory, unified memory and its page-fault cost
• Global-memory coalescing: 32-byte sectors, AoS vs SoA, the strided-access cliff
• Shared memory: __shared__, tiling, bank conflicts, __syncthreads() discipline
• Occupancy in practice: registers/thread, __launch_bounds__, the occupancy calculator
• Streams and events: async copies, overlap, cudaEventElapsedTime for honest timing
• Atomics and reductions: warp shuffle (__shfl_down_sync), block reduction patterns
• Profiling: Nsight Systems (timeline) vs Nsight Compute (kernel deep-dive); the three metrics that matter first

Labs

Lab 01 — First Kernels: Vector Add → Tiled Matmul (CUDA C++)

Lab 02 — Coalescing, Occupancy & the Profiler (CUDA C++)

Deliverables Checklist

• Both labs build and pass their built-in verification on a real GPU
• Timing table: naive vs tiled matmul, with the roofline explanation written down
• Bandwidth tables for stride / AoS-SoA / transpose experiments
• PROFILE.md filled: each kernel's limiter named with ncu evidence
• You can write the CUDA_CHECK macro and the launch-config formula from memory

Interview Relevance

• "Launch config for a 1920×1080 image filter — walk me through the math."
• "Your kernel gets 8% of peak bandwidth. List your first three hypotheses."
• "Why are thread blocks required to be independent? What feature does that enable?"
• "What does __syncthreads() do, and what happens if it's inside a divergent branch?"
• "When is unified memory acceptable in production?"

Warmup Guide — CUDA Programming

Zero-to-expert primer for Phase 02. Assumes Phase 01's machine model (warps, SMs, memory hierarchy, roofline) and basic C/C++. By the end you can write, verify, time, and profile real kernels — and review other people's.

Table of Contents

• Chapter 1: The CUDA Stack — What You're Actually Calling
• Chapter 2: The Programming Model — Grid, Block, Thread
• Chapter 3: Your First Kernel, Line by Line
• Chapter 4: Error Handling — The Discipline That Separates Professionals
• Chapter 5: The Memory API — cudaMalloc to Unified Memory
• Chapter 6: Coalescing — The #1 Performance Rule
• Chapter 7: Shared Memory and Tiling — Raising Arithmetic Intensity
• Chapter 8: Occupancy in Practice
• Chapter 9: Streams, Events, and Honest Timing
• Chapter 10: Reductions and Warp Primitives
• Chapter 11: Profiling — Nsight Systems vs Nsight Compute
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: The CUDA Stack — What You're Actually Calling

Zero background: "CUDA" is three things people conflate: a language extension (C++ plus __global__, <<<>>>), a runtime library (libcudart — the cudaMalloc/cudaMemcpy API), and a driver stack (libcuda.so, the kernel-mode driver — Phase 04 territory).

The call chain when you launch a kernel:

your code → CUDA Runtime API (libcudart)   cudaLaunchKernel(...)
          → CUDA Driver API   (libcuda.so) cuLaunchKernel(...)
          → ioctl(2) into the kernel-mode driver (/dev/nvidia*)
          → commands written to a ring buffer the GPU hardware consumes

Two facts with production consequences:

1. The runtime API is a convenience layer over the driver API. Inference engines and runtimes (Phase 05/09) often use the driver API directly for explicit context control, loading cubins at runtime (cuModuleLoad), and multi-tenancy. You should be able to read both.
2. Kernel launches are asynchronous. The CPU enqueues and returns in ~5–10 µs; the GPU executes later. Everything about timing (Ch. 9), error reporting (Ch. 4), and overlap (streams) follows from this one fact.

Misconception: "nvcc is a compiler like gcc." nvcc is a driver that splits your file: host code → your host compiler; device code → PTX/SASS through NVIDIA's backend. Phase 03 dissects this pipeline; here you just use it.

Chapter 2: The Programming Model — Grid, Block, Thread

You write one scalar function (a kernel); the runtime launches N instances:

• Thread: one instance. Knows its coordinates via threadIdx, blockIdx, blockDim, gridDim (each a 3-component vector for 1D/2D/3D convenience).
• Block: up to 1,024 threads that share a shared-memory allocation, can __syncthreads(), and are guaranteed to run on one SM.
• Grid: all blocks of a launch. Blocks must be independent — no cross-block sync within a kernel, no assumption about execution order.

Why block independence is the load-bearing rule: it's what lets the hardware scheduler scatter blocks across however many SMs exist — 10 on a laptop, 132 on H100 — and retire/replace them in any order. Block independence is CUDA's scalability contract; it is why 2008 kernels still scale on 2025 silicon. (It's also why a "wait for block X" spinlock in a kernel deadlocks: block X may not be resident.) Cooperative groups (cudaLaunchCooperativeKernel) exist as the explicit, occupancy-checked exception.

The index formula you'll write a thousand times:

int i = blockIdx.x * blockDim.x + threadIdx.x;     // global index
if (i < n) { ... }                                  // guard the tail
int nblocks = (n + threads - 1) / threads;          // ceil-div launch config

The guard exists because the grid is a multiple of the block size and n usually isn't. Forgetting it = out-of-bounds writes = the classic first CUDA bug (intermittent corruption, not a crash — see Ch. 4).

Mapping back to Phase 01: blocks are carved into warps (threads 0–31, 32–63, …). All divergence/coalescing reasoning operates at warp granularity within the block you configured. Block size should essentially always be a multiple of 32 — 128 or 256 are the standard defaults.

Chapter 3: Your First Kernel, Line by Line

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}
// host:
int threads = 256;
int blocks  = (n + threads - 1) / threads;
vec_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

• __global__: callable from host, runs on device. (__device__ = device-only helper; __host__ __device__ = compiled for both.)
• The pointers are device pointers (from cudaMalloc). Dereferencing a host pointer in a kernel is undefined behavior — typically illegal memory access reported much later (Ch. 4).
• <<<blocks, threads>>> has two more slots: dynamic shared-memory bytes and a stream — <<<blocks, threads, smem_bytes, stream>>> (Ch. 7, 9).

Production reality: vector add is the memory-bound kernel (AI ≈ 0.08, Phase 01 Ch. 7). Its only performance question is "did we hit memory bandwidth?" — which makes it the perfect coalescing test harness, and exactly how Lab 01 uses it.

Chapter 4: Error Handling — The Discipline That Separates Professionals

CUDA errors come in two flavors, and conflating them wastes engineer-days:

• Synchronous: bad launch config, OOM at cudaMalloc — returned immediately by the call.
• Asynchronous: the kernel itself faults (out-of-bounds, bad pointer) — surfaces at the next synchronizing call, attributed to an innocent line.

The non-negotiable macro (write it from memory in interviews):

#define CUDA_CHECK(call) do {                                         \
    cudaError_t e = (call);                                           \
    if (e != cudaSuccess) {                                           \
        fprintf(stderr, "CUDA error %s:%d: %s\n",                     \
                __FILE__, __LINE__, cudaGetErrorString(e));           \
        exit(1);                                                      \
    }                                                                 \
} while (0)

kernel<<<g, b>>>(args);
CUDA_CHECK(cudaGetLastError());        // catches launch-config errors
CUDA_CHECK(cudaDeviceSynchronize());   // catches kernel-execution errors (debug builds)

Debugging ladder when something faults: 1) compute-sanitizer ./app (catches the faulting kernel + address — the GPU's valgrind), 2) CUDA_LAUNCH_BLOCKING=1 (makes async errors synchronous so attribution is correct), 3) cuda-gdb for the rare hard case.

Production significance: a fault in any kernel poisons the CUDA context — subsequent calls return errors until process restart. This is why serving engines treat "CUDA error" as fatal-and-restart (Phase 10 runbooks), and why your platform must isolate tenants at process granularity, not context granularity (Phase 06).

Chapter 5: The Memory API — cudaMalloc to Unified Memory

The four tiers, in order of how often you should reach for them:

1. Explicit — cudaMalloc + cudaMemcpy. Predictable, fast, verbose. The production default. cudaMemcpy is synchronous-ish; the async variant requires tier 2.
2. Pinned host memory — cudaMallocHost. Page-locked so the GPU's DMA engine can read it directly: ~2× faster transfers and the prerequisite for cudaMemcpyAsync actually being async (pageable copies silently stage through a pinned bounce buffer and serialize). Cost: pinned pages are stolen from the OS; pin gigabytes, not tens of gigabytes.
3. Unified memory — cudaMallocManaged. One pointer valid on both sides; pages migrate on demand. Beautiful for prototypes; in production the page-fault storms on first touch and on every CPU↔GPU ping-pong are a classic mystery slowdown (visible as Unified Memory rows in Nsight Systems). cudaMemPrefetchAsync is the remedy when you must use it.
4. Stream-ordered pools — cudaMallocAsync/cudaFreeAsync. Allocation as a stream operation from a driver-managed pool — because raw cudaMalloc/Free synchronize the device and fragment, which is also why every framework built its own caching allocator. Phase 05 Lab 01 has you build exactly that allocator in Rust — this chapter is its motivation.

The rule that survives all four tiers: minimize transfers; batch them; overlap them with compute (Ch. 9). PCIe is 50× narrower than HBM (Phase 01 Ch. 6) — a synchronous host↔device copy inside a hot loop is the single most common production CUDA performance bug.

Chapter 6: Coalescing — The #1 Performance Rule

The mechanism: when a warp executes a load, the hardware takes all 32 addresses and computes which 32-byte sectors of memory they touch, then issues one transaction per distinct sector.

• 32 consecutive floats → 4 sectors → minimal transactions. Coalesced.
• Stride-32 floats (e.g., walking a column of a row-major matrix) → 32 different sectors → 32 transactions, each delivering 4 useful bytes of 32 fetched. ~8× wasted bandwidth, and worse with bigger strides.

This is the GPU's version of Phase 01 Lab 02's cache-line stride cliff — same physics, warp-granular.

The three patterns to know:

1. Adjacent threads ↔ adjacent addresses. a[blockIdx.x*blockDim.x + threadIdx.x] is perfect; a[threadIdx.x * pitch] is the anti-pattern.
2. SoA beats AoS. struct {float x,y,z;} pts[n] makes pts[i].x loads stride-12; three arrays x[n], y[n], z[n] coalesce each field. This is why GPU-facing data layouts (and Phase 07's KV-cache layouts) are structure-of-arrays.
3. Matrix transpose is the canonical exercise: naive transpose must uncoalesce either reads or writes; the fix is staging a tile through shared memory so both global phases are coalesced — Lab 02 implements it, padding included (Ch. 7).

Misconception: "the L2 cache makes coalescing obsolete." L2 softens repeated sloppiness, but a bandwidth-bound kernel wasting 8× on sectors is still ~8× slow — Lab 02's table is the proof you'll carry around.

Chapter 7: Shared Memory and Tiling — Raising Arithmetic Intensity

Phase 01 Ch. 7 said: performance = min(peak, AI × bandwidth). Tiling is the technique for raising AI, and shared memory is its vehicle.

Naive matmul: each thread computes one C[i][j], reading a full row of A and column of B from global memory → every input element fetched N times → AI ≈ O(1) → bandwidth-bound at a few % of peak FLOPs.

Tiled matmul (Lab 01's centerpiece):

__shared__ float As[T][T], Bs[T][T];
for (int t = 0; t < N/T; t++) {
    As[ty][tx] = A[row*N + t*T + tx];      // coalesced cooperative load
    Bs[ty][tx] = B[(t*T + ty)*N + col];    // coalesced cooperative load
    __syncthreads();                        // tile fully loaded before use
    for (int k = 0; k < T; k++) acc += As[ty][k] * Bs[k][tx];
    __syncthreads();                        // all done reading before overwrite
}

Each global element is now loaded once per tile instead of once per thread → AI multiplied by T (tile dim, typically 16–32) → 5–10× speedup, exactly as the roofline predicts. cuBLAS continues the same idea further down the hierarchy: register-blocking (each thread computes a small output micro-tile from registers) and Tensor Core fragments — that's the remaining gap when Lab 01's extension compares against sgemm.

The two rules of __syncthreads(): (1) it's a block-wide barrier — every thread of the block must reach it, so it must never sit inside a branch that's divergent across the block (deadlock/UB); (2) you need it twice per tile — after loading (before consuming) and after consuming (before the next overwrite). Lab 01 asks you to remove one and observe the corruption.

Bank conflicts: shared memory has 32 banks, word-interleaved. tile[threadIdx.x][k] with a 32-wide tile puts a warp's accesses in one bank (32-way conflict → serialized). Fix: pad the tile — __shared__ float tile[32][33] — the +1 rotates each row's bank alignment. Lab 02 measures before/after; it's a ~30–40% transpose-bandwidth swing for one character of code.

Chapter 8: Occupancy in Practice

Phase 01 Ch. 5 defined occupancy as a latency-hiding budget. The practice:

What limits it (per SM, Ampere numbers): 65,536 registers, 100–164 KB shared memory, 2,048 thread slots, 32 block slots. Whichever resource your kernel exhausts first caps how many blocks are resident:

• 256 threads/block, 64 regs/thread → 16,384 regs/block → 4 blocks → 1,024 threads = 50% occupancy (register-limited).
• 48 KB smem/block on a 100 KB SM → 2 blocks resident regardless of registers.

The tools: nvcc -Xptxas=-v prints regs/smem per kernel; cudaOccupancyMaxActiveBlocksPerMultiprocessor computes it at runtime (use it to pick launch configs in library code — Phase 09's HAL does this); __launch_bounds__(maxThreads, minBlocks) tells the compiler to fit a target (it will spill registers to comply — measure, don't assume).

The judgment call (this is what reviewers actually argue about): more occupancy hides more latency, but fewer registers per thread can mean spills, and less shared memory per block can mean less tiling/reuse. There is no universally right answer — only the profile. Volkov again: high-ILP kernels can peak at 25% occupancy. The professional posture: state the limiter, show the ncu evidence, justify the chosen point.

Chapter 9: Streams, Events, and Honest Timing

Streams are ordered work queues; operations in different streams may overlap. The legacy default stream synchronizes with everything (compile with --default-stream per-thread or just use explicit streams everywhere — the platform-code default).

The canonical overlap pattern — pipelining chunked transfers against compute (this is also Phase 07's prefill/decode overlap in miniature):

for (int c = 0; c < nchunks; c++) {
    cudaMemcpyAsync(d_in + off, h_in + off, bytes, H2D, stream[c % 2]);
    kernel<<<g, b, 0, stream[c % 2]>>>(d_in + off, d_out + off);
    cudaMemcpyAsync(h_out + off, d_out + off, bytes, D2H, stream[c % 2]);
}

Requirements for actual overlap: pinned host memory (Ch. 5), separate streams, and chunk sizes big enough to amortize launch overhead (~10 µs/launch — relevant to Phase 05's scheduler and why CUDA Graphs exist for short-kernel pipelines).

Honest timing — the three rules, violated in 90% of bad benchmark claims:

1. Time with events (cudaEventRecord/cudaEventElapsedTime), or wall-clock only around an explicit cudaDeviceSynchronize(). Timing an async launch with std::chrono measures the enqueue, not the kernel.
2. Warm up first (first launch includes module load/JIT — Phase 03 explains the JIT path).
3. Report the median of many reps, with clocks noted if comparing across sessions (GPUs thermal-throttle; nvidia-smi -q -d CLOCK).

Events also serve as cross-stream dependencies (cudaStreamWaitEvent) — the primitive from which Phase 05 Lab 02's stream scheduler builds its DAG.

Chapter 10: Reductions and Warp Primitives

Summing N elements exposes everything this phase taught at once — and it's a top-3 interview exercise.

The evolution every CUDA engineer should be able to narrate:

1. Atomic per element: atomicAdd(&out, a[i]) — correct, serializes on one address, terrible.
2. Shared-memory tree: block loads to smem, halving strides with __syncthreads() between levels, thread 0 atomically adds the block's sum. Classic; bank-conflict-aware versions use sequential addressing (stride from blockDim/2 down, not interleaved by 2× — the interleaved version also has divergence problems: if (tid % (2*s) == 0) is the per-lane branch Phase 01 warned about).
3. Warp shuffle finish: the last 32 elements don't need shared memory at all — val += __shfl_down_sync(0xffffffff, val, offset) moves data register-to-register within the warp, no smem, no sync. Modern reductions: shuffle within warps, one smem round between warps, shuffle again.
4. (Library answer: CUB BlockReduce — and saying "in production I'd use CUB, here's what it does inside" is the strongest possible interview answer.)

Warp primitives to know by name: __shfl_*_sync (data exchange), __ballot_sync / __any_sync / __all_sync (votes — Lab 01 of Phase 01 simulated vote.any), __activemask(). All take the mask-first argument because of Volta independent thread scheduling (Phase 01 Ch. 3).

Chapter 11: Profiling — Nsight Systems vs Nsight Compute

Two tools, two altitudes — using the wrong one wastes afternoons:

• Nsight Systems (nsys) — the timeline. CPU threads, API calls, transfers, kernel spans, stream overlap. Answers: "where does wall-clock time go?" "is the GPU idle between kernels?" "did my copies overlap?" Always start here. nsys profile -o report ./app → open in the GUI or nsys stats.
• Nsight Compute (ncu) — the microscope. Per-kernel hardware counters: achieved bandwidth, warp execution efficiency (Lab 01 of Phase 01 computed this by hand!), occupancy achieved vs theoretical, stall reasons, and a built-in roofline chart. ncu --set full ./app. Expensive (replays kernels many times).

The first-pass triage (memorize as a flowchart):

1. nsys: GPU mostly idle? → it's a pipeline problem (launch overhead, sync points, transfers) — no kernel tuning will help.
2. GPU busy → ncu the top kernel: SOL Memory% vs SOL Compute% tells you which roof you're under (Phase 01's roofline, now measured for you).
3. Memory-bound → check gld_efficiency/sectors-per-request (coalescing, Ch. 6), then data reuse (tiling, Ch. 7). Compute-bound → precision/Tensor-Core utilization. Neither high but stalls high → latency-bound: occupancy (Ch. 8) or dependency chains.

This triage is Lab 02's PROFILE.md exercise, and — at the leadership level — it's the checklist you require attached to any "we optimized the kernel" PR.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02. Lab 01 builds the kernels; Lab 02 measures their sins.

Lab 01 (first kernels):

1. Build and run first (make && ./kernels); confirm all PASS lines, then read kernels.cu against Chapters 2–4 and 7.
2. Work the verification path: break the tail guard deliberately, see the corruption, fix it, then break a __syncthreads() and see the race.
3. Extensions in order: rectangular tiles → register micro-tiles → cuBLAS comparison (expect cuBLAS 2–5× ahead; explain why using Ch. 7's last paragraph).
4. No local GPU: the README's Colab cell runs everything; nvcc flags identical.

Lab 02 (coalescing & profiler):

1. Run the benchmark and fill the tables before profiling — predict, then verify with ncu (prediction-first is the habit that makes profiling stick).
2. The transpose ladder (naive → tiled → padded) is the heart: tie each step to Ch. 6/7 mechanisms in one sentence each in PROFILE.md.
3. Finish with the triage flowchart (Ch. 11) applied to both your matmuls from Lab 01 — which roof is each under, with evidence.

Success Criteria

• You can write a correct kernel + launch + CUDA_CHECK discipline from a blank editor with no references (Ch. 2–4)
• You can explain sync vs async errors and run the 3-step debugging ladder (Ch. 4)
• You can state the coalescing mechanism in terms of warp addresses → 32-byte sectors, and the SoA and transpose consequences (Ch. 6)
• Your tiled matmul beats naive ~5–10× and you can derive why from the roofline, including both __syncthreads() placements (Ch. 7)
• You can name a kernel's occupancy limiter from -Xptxas=-v output (Ch. 8)
• Your benchmark numbers come from events with warmup and reps (Ch. 9)
• You can narrate the reduction evolution through warp shuffles (Ch. 10)
• Given an nsys+ncu pair, you can name the limiter and the next experiment in under two minutes (Ch. 11)

Interview Q&A

Q1: Walk me through launching a blur filter over a 1920×1080 image. A: Pick a 2D block, say 16×16 = 256 threads. Grid = ceil(1920/16) × ceil(1080/16) = 120 × 68 blocks. Each thread computes x = blockIdx.x*16 + threadIdx.x, y likewise, guards if (x < 1920 && y < 1080) because 1080/16 = 67.5 rounds up and the last block row hangs off the edge. For a blur, neighboring threads read overlapping pixels, so I'd stage a (16+2r)² halo tile in shared memory with cooperative coalesced loads, __syncthreads(), then filter from smem. Launch asynchronously into a stream, CUDA_CHECK(cudaGetLastError()) after.

Q2: Your kernel achieves 8% of peak bandwidth. First three hypotheses? A: (1) Uncoalesced access — warp addresses scattering across 32-byte sectors: check ncu sectors-per-request (~4 is ideal for floats; 32 means fully strided) — typical causes: column-major walk, AoS layout, bad pitch. (2) Latency-bound, not bandwidth-bound — too few warps in flight to saturate the memory system: check achieved occupancy and stall reasons; fix with more parallelism per launch or more ILP per thread. (3) The kernel isn't actually the time sink — short kernel dominated by launch overhead or serialized with transfers: check the nsys timeline first. (Bonus: unified-memory page faults masquerading as kernel time.)

Q3: Why must thread blocks be independent? A: Independence is the scalability contract: the hardware can place, execute, and retire blocks in any order on any SM, which lets one binary saturate a 10-SM laptop and a 132-SM H100, and lets the scheduler backfill SMs as blocks finish (Phase 01's latency-hiding at block granularity). Inter-block sync within a kernel would require all blocks resident simultaneously — deadlock otherwise. The sanctioned escape hatches: end the kernel (grid-wide barrier = kernel boundary), atomics for unordered communication, or cooperative launch, which checks co-residency at launch time.

Q4: What goes wrong with __syncthreads() in divergent code? A: It's a block-wide barrier; every thread must execute the same barrier. If some threads of a block skip it (a branch divergent at block scope), the arriving threads wait forever — deadlock or undefined behavior, in practice a hang or corrupted results. Warp-level divergence reconverges by hardware, but __syncthreads() semantics demand all-threads participation: the rule is no __syncthreads() under conditions that vary within the block; restructure so the barrier is unconditional (compute the condition, sync, then branch).

Q5: Pinned vs pageable vs unified memory — when each? A: Pageable (malloc) host memory forces the driver to stage transfers through an internal pinned bounce buffer — slower, and cudaMemcpyAsync silently degrades to effectively synchronous. Pinned (cudaMallocHost) gives DMA-direct ~2× transfer speed and true async — required for any overlap pipeline; cost is stealing unpageable RAM from the OS, so pin transfer buffers, not datasets. Unified (cudaMallocManaged) trades control for convenience: on-demand page migration is fine for prototypes and oversubscription-with-prefetch on data too big for HBM, but fault storms on alternating CPU/GPU touches are a notorious production slowdown — in serving infrastructure I require explicit transfers at review time unless there's a measured case.

Q6 (leadership): An engineer's PR claims a 3× kernel speedup. What do you require before merge? A: (1) Correctness evidence: elementwise comparison against a reference within a stated tolerance, plus a compute-sanitizer clean run — fast-but-wrong kernels are the classic regression. (2) Honest methodology: event-based timing, warmup, median of N, fixed clocks noted, same input distribution as production. (3) The mechanism: which limiter moved, with before/after ncu evidence (e.g., sectors-per-request 32→4, SOL Memory 85%) — "I don't merge speedups we can't explain, because unexplained wins are usually measurement bugs." (4) Portability notes: which architectures it was tuned on, occupancy impact at other block sizes (Phase 09's HAL concern). This checklist is Lab 02's PROFILE.md, generalized into team process.

References

• NVIDIA, CUDA C++ Programming Guide — https://docs.nvidia.com/cuda/cuda-c-programming-guide/
• NVIDIA, CUDA C++ Best Practices Guide — https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
• Mark Harris, "Optimizing Parallel Reduction in CUDA" (the classic 7-step deck) — https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
• Mark Harris, "An Efficient Matrix Transpose in CUDA C/C++" — https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/
• NVIDIA Nsight Systems / Nsight Compute documentation — https://docs.nvidia.com/nsight-systems/ , https://docs.nvidia.com/nsight-compute/
• CUTLASS (how production GEMMs are actually structured) — https://github.com/NVIDIA/cutlass
• Sanders & Kandrot, CUDA by Example (dated but the gentlest on-ramp)
• Cheng, Grossman, McKercher, Professional CUDA C Programming
• Cross-track: Phase 01 WARMUP (machine model), Phase 03 WARMUP (what nvcc does with this code)

🛸 Hitchhiker's Guide — Phase 02: CUDA Programming

Read this if: you understand the Phase 01 machine model but haven't shipped CUDA. This is the field guide: idioms, numbers, foot-guns, and the review standards a GPU platform team should hold.

0. The 30-second mental model

CUDA = C++ where one function runs N times concurrently, N = grid × block. You manage three things: indices (which element am I?), memory (is my warp's access coalesced? is reuse staged in shared memory?), and asynchrony (launches return immediately; streams order work; events time and connect it). Performance work = move the kernel toward the correct roof from Phase 01's roofline, then stop.

1. The idioms you'll type forever

// ceil-div launch + tail guard — THE pattern
int t = 256, b = (n + t - 1) / t;
k<<<b, t>>>(d_x, n);                 // d_x from cudaMalloc, never a host ptr
// error discipline
CUDA_CHECK(cudaGetLastError());      // launch errors (sync)
CUDA_CHECK(cudaDeviceSynchronize()); // kernel errors (async) — debug builds
// honest timing
cudaEventRecord(t0); k<<<...>>>(...); cudaEventRecord(t1);
cudaEventSynchronize(t1); cudaEventElapsedTime(&ms, t0, t1);

Block size: multiple of 32, default 128/256, decided by measurement not vibes. 3D grids are indexing sugar — the hardware sees a flat list of blocks.

2. Memory decisions, ranked by frequency of regret

1. Sync copies in hot loops (PCIe is 1/50th of HBM): batch, pin, overlap.
2. Uncoalesced layouts: AoS structs, column walks of row-major data, threadIdx.y-fastest indexing. Fix = SoA + adjacent-thread-adjacent-address. ncu smoking gun: sectors/request ≈ 32 (want ~4).
3. Unified memory in serving paths: page-fault storms. Managed memory is a prototyping tool; production transfers are explicit.
4. Raw cudaMalloc/Free per request: device-synchronizing and fragmenting — use cudaMallocAsync pools or a caching allocator (you'll build one in Phase 05; PyTorch's is the reference design).
5. Forgetting pinned memory for async: cudaMemcpyAsync from pageable memory quietly serializes.

3. Shared memory: the two-line contract

• Tile global→smem with coalesced cooperative loads; __syncthreads(); compute from smem; __syncthreads(); next tile. Matmul AI goes from O(1) to O(T): the whole game in one sentence.
• Pad [32][33] to dodge bank conflicts; never put __syncthreads() under a block-divergent branch.

4. Numbers worth carrying

5. The profiler triage (commit to memory)

nsys timeline:
  GPU idle?           → pipeline problem: launches, syncs, transfers. Stop kernel-tuning.
  GPU busy → ncu the top kernel:
    SOL Memory high?  → coalescing → reuse/tiling → vectorize loads (float4)
    SOL Compute high? → precision, Tensor Cores, better algorithm
    both low?         → latency-bound: occupancy, ILP, dependency chains

The strongest habit this phase teaches: predict before measuring. Write down the expected bandwidth/speedup from the roofline, then check. A team that predicts-then-profiles compounds; a team that profiles-then-rationalizes doesn't.

6. Debugging war stories (so you recognize them)

• "It crashes on a line that's obviously fine" — async error from an earlier kernel. CUDA_LAUNCH_BLOCKING=1, find the real culprit.
• "Works at n=1000, corrupts at n=1000000" — missing tail guard; or grid dim overflowed 65,535 in a .y/.z dimension.
• "Slower after I 'optimized' registers" — -maxrregcount caused spills; occupancy went up, performance went down. Volkov was right.
• "First call takes 200 ms" — context init + module JIT. Warm up; ship SASS for your target arch (Phase 03's fatbin story).
• "Results differ run to run by 1e-6" — float reduction order changes with block scheduling; that's not a bug, it's floating-point non-associativity. Define tolerance-based verification (and lock seeds for argmax-style outputs).
• "Everything returns errors after one bad kernel" — poisoned context; restart the process. Architect tenant isolation accordingly (Phase 06).

7. From engineer to head-of-engineering

What changes at your level isn't the syntax — it's the standards:

• Kernel PR checklist: reference-comparison test, sanitizer-clean, event-based timing with reps, ncu before/after naming the limiter, arch notes. (Lab 02 builds it; your job is making it stick.)
• Estimate sanity: when someone says "we'll get 10× by porting to GPU," you now ask: what's the AI? what fraction is parallel? how much PCIe traffic? Amdahl + roofline kills most 10× claims in five minutes — kindly, with math.
• Hiring signal: ask candidates the reduction-evolution question (§Ch. 10 of the WARMUP). Depth on "why shuffle beats shared memory for the last warp" separates practitioners from tutorial-readers.

Next: Phase 03 — what nvcc actually does to this code (PTX, SASS, and a mini compiler you'll write yourself), because a platform that promises "independence from underlying hardware" lives or dies in the compiler layer.

Lab 01 — First Kernels: Vector Add → Tiled Matmul

Phase: 02 — CUDA Programming | Difficulty: ⭐⭐⭐☆☆ | Time: 4–6 hours Language: CUDA C++ | Hardware: any NVIDIA GPU (sm_70+); Colab T4 works

Concept primer: ../WARMUP.md Ch. 2–4, 7, 9.

Run

make            # nvcc -O3 -arch=native (or ARCH=sm_75 make for Colab T4)
./kernels

No local GPU? Google Colab

Runtime → Change runtime type → T4 GPU, then in one cell:

!nvcc --version
%%writefile kernels.cu
<paste kernels.cu>
!nvcc -O3 -arch=sm_75 -o kernels kernels.cu && ./kernels

0. The mission

Three kernels, one discipline:

1. vec_add — the programming model + error handling + the tail guard.
2. matmul_naive — one thread per output element, all reads from global memory.
3. matmul_tiled — the same math with shared-memory tiling.

Every kernel is verified elementwise against a CPU reference, and timed with CUDA events (warmup + median of reps). Expected result: tiled beats naive ~5–10×, because tiling raises arithmetic intensity by the tile factor — the Phase 01 roofline argument made executable.

Sample output (T4, n=1024):

vec_add        : PASS   33.55 GB/s effective (memory-bound, as designed)
matmul_naive   : PASS   612 ms   0.35 TFLOP/s
matmul_tiled   : PASS    78 ms   2.75 TFLOP/s   speedup 7.8x  [PASS >= 4x]

1. Read the code in this order

1. CUDA_CHECK and check_close — the discipline. Note where cudaGetLastError is called (right after launch — catches config errors) vs cudaDeviceSynchronize (catches the kernel's own faults).
2. vec_add — the index formula and tail guard. Exercise: delete the if (i < n) guard, run with n=1<<20+3, observe either corruption or a sanitizer report (compute-sanitizer ./kernels). Put it back.
3. matmul_naive — each thread reads a row of A and a column of B from global memory. The column read of B is actually coalesced across threads (adjacent tx → adjacent B columns at fixed k) — the problem isn't coalescing here, it's zero reuse: every element fetched N times. AI ≈ O(1).
4. matmul_tiled — the cooperative load / sync / compute / sync rhythm (WARMUP Ch. 7). Exercise: comment out the second __syncthreads() and run — verification fails intermittently (a fast warp overwrites the tile while a slow one still reads). This is the cheapest race-condition education available; take it.
5. The timing harness — events, warmup, median-of-5. Compare against the naive std::chrono around an async launch (also included, prints the lie).

2. Why tiled wins — write this down in your own words

Naive: each of N³ multiply-adds needs 2 fresh global loads → AI ≈ 1/8 FLOP/byte → the bandwidth roof at ~320 GB/s (T4) caps you near 0.04 TFLOP/s per... except L2 catches some reuse, so you observe ~0.3. Tiled with T=32: each global element is loaded once per tile and reused 32 times from shared memory → AI × 32 → the roof rises past 1 TFLOP/s on the same silicon. Same FLOPs, same GPU, 8× — layout and staging, not arithmetic.

The remaining 3–5× to cuBLAS: register micro-tiles (each thread computes 4–8 outputs from registers, raising AI again), double-buffered async tile loads (cp.async), Tensor Cores via mma, and per-arch tile-shape tuning. That ladder is CUTLASS's table of contents — and the extension exercises walk its first rung.

3. Extension ladder

1. Rectangular blocking: 64×16 tiles; measure. Why might non-square help? (Hint: B's tile loads per output column.)
2. Register micro-tile: each thread computes a 2×2 patch of C. Expect another 1.5–2.5×. You're now register-blocked like real GEMMs.
3. cuBLAS comparison: link -lcublas, time cublasSgemm, report your % of it. >30% with micro-tiles is respectable; explain the rest.
4. __launch_bounds__(256, 4) on the tiled kernel: check -Xptxas=-v register count before/after, occupancy, and time. Did the compiler spill?

4. Common pitfalls

1. Host pointer passed to a kernel — illegal access at the next sync, blamed on the wrong line (CUDA_LAUNCH_BLOCKING=1 to localize).
2. Verifying with == on floats — reduction/FMA order differs from CPU; use relative tolerance (the harness uses 1e-3 relative).
3. Timing the enqueue, not the kernel (chrono without sync) — the harness prints both so you see the discrepancy once and never trust chrono again.
4. -arch mismatch — sm_90 binary on a T4 dies at load; -arch=native locally, explicit ARCH= on shared machines (the full fatbin story is Phase 03).

5. What this lab proves about you

You can produce a verified, honestly-benchmarked kernel and explain its performance from the roofline — the bar you will later hold every kernel PR to (Lab 02 turns that into a checklist).

Lab 02 — Coalescing, Bank Conflicts & the Profiler

Phase: 02 — CUDA Programming | Difficulty: ⭐⭐⭐⭐☆ | Time: 4–6 hours Language: CUDA C++ + Nsight Compute CLI | Hardware: any NVIDIA GPU; Colab T4 works

Concept primer: ../WARMUP.md Ch. 6–8, 11.

Run

make && ./bench               # bandwidth tables: stride, AoS/SoA, transpose ladder
ncu --set basic ./bench       # then: the profiler exercise (see §3)

0. The mission

Three experiments that turn WARMUP Chapter 6–7 claims into numbers you measured, then a profiler session that teaches you to name limiters with evidence:

1. Stride sweep: copy bandwidth at stride 1, 2, 4, 8, 16, 32 — watch coalescing die.
2. AoS vs SoA: load one field from a 3-float struct array vs a flat array.
3. Transpose ladder: naive → shared-memory tiled → tiled+padded; recover copy bandwidth step by step.

Expected shape of results (T4; absolute numbers vary):

== stride sweep (effective GB/s) ==
stride  1: 240.1     stride  4:  68.3     stride 16:  17.9
stride  2: 124.5     stride  8:  34.6     stride 32:   9.2     <- ~26x cliff

== AoS vs SoA (loading .x only) ==
AoS:  81.4 GB/s   SoA: 238.0 GB/s   (~3x: the other 8 bytes ride the sectors)

== transpose ladder ==
copy (ceiling) : 235 GB/s
naive          :  62 GB/s   (writes uncoalesced)
tiled          : 148 GB/s   (both phases coalesced; bank conflicts on read)
tiled + pad    : 221 GB/s   (94% of copy — done)

1. Why each step behaves as it does

• Stride s: each warp load touches ~min(32, s·4·32/32) distinct 32-byte sectors; useful bytes per sector fall as 1/s until one float per sector (s ≥ 8) — after which the cliff flattens: you're paying one full sector per element. The same curve you produced on the CPU in Phase 01 Lab 02, now at warp granularity.
• AoS: reading pts[i].x drags .y and .z through the memory system as sector ballast. SoA puts 32 consecutive .x in 4 sectors. (This single diagram is why every GPU-resident data structure in Phases 05–07 is SoA.)
• Transpose naive: reads coalesce (row-major in), writes scatter (column-major out) — one direction always loses. Tiled: stage a 32×32 tile in shared memory; read coalesced, write coalesced, the "turn" happens in smem. Padding [32][33]: the tiled version's smem reads hit one bank 32 ways; the +1 column shifts each row's bank phase → conflict-free. One character, ~1.5× on the transpose.

2. Record your tables

Fill in the actual numbers in PROFILE.md §1–3. Predict each before running — prediction-first is the habit.

3. The profiler session (§4 of PROFILE.md)

On the naive and padded transpose kernels:

ncu --kernel-name regex:transpose --set basic ./bench

Find and record, for each kernel:

Then write the verdict line for each kernel in PROFILE.md: "<kernel> is <memory/compute/latency>-bound because <two metrics>; the next experiment would be <X>." — that sentence format is the whole point of the lab. It's also the PR-review standard you'll set in Phase 12.

On Colab: !ncu --set basic ./bench works on T4 instances (driver permitting; if ERR_NVGPUCTRPERM, add --target-processes all or use nsys profile --stats=true ./bench for the timeline-level view instead and note the difference in what you can conclude).

4. Extensions

• float4 vectorized copy: 16-byte loads per thread; measure against stride-1 float copy. Why does it help? (Fewer instructions per byte; wider sectors per request.)
• cp.async (sm_80+): async global→shared copies for the tiled transpose; overlap the next tile's load with this tile's store.
• Occupancy experiment: add __launch_bounds__(1024, 1) to the padded transpose, check -Xptxas=-v and achieved occupancy, explain the change.
• Write the checklist: distill §3 into a one-page "kernel PR review checklist" — you will reuse it verbatim in Phase 12's leadership artifacts.

5. Common pitfalls

1. Comparing kernels at different clocks (thermal throttle mid-run): lock with nvidia-smi -lgc where permitted, or interleave reps.
2. Profiling with --set full on a long benchmark — replay overhead makes it crawl; use basic + targeted metrics.
3. Concluding "memory-bound" from low compute% alone — check both SOL numbers; both-low means latency-bound, a different fix.
4. Forgetting that ncu serializes kernels — overlap effects you saw in nsys vanish under ncu; the tools answer different questions (WARMUP Ch. 11).

6. What this lab proves about you

You can quantify the two memory sins (uncoalesced access, bank conflicts), fix them with the two standard remedies (SoA/tiling, padding), and defend the analysis with profiler evidence — the exact competency a GPU platform lead needs to review kernel work credibly.

Phase 03 — GPU Compilers: NVCC, PTX, SASS, and Codegen

Difficulty: ⭐⭐⭐⭐☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], GPU Compiler Engineer, Runtime Engineer Hardware needed: none for Lab 02 (the compiler you build emits and interprets PTX); any NVIDIA GPU (or Godbolt) for Lab 01

Why This Phase Exists

A platform that "operates independently of the underlying hardware" (the JD's first bullet) is, at its core, a compiler and runtime story: something must translate a hardware-neutral program representation into each vendor's ISA, and do it well enough that customers don't pay a performance tax for portability. Every player in this market — CUDA itself (PTX as the stable virtual ISA), Triton, MLIR, ROCm, OpenAI/vLLM kernels, every startup accelerator's "we support PyTorch" claim — is a position in the compiler argument.

You cannot lead that platform if the compiler layer is a black box. This phase opens it: you will trace nvcc's real pipeline, read PTX and SASS with comprehension, understand JIT vs AOT and the fatbin compatibility model (a production operations concern, not trivia), and then write a working compiler that lexes, parses, and emits legal PTX — plus an interpreter that executes your PTX per-thread so you can verify semantics without a GPU.

Concepts

• The nvcc pipeline end-to-end: host/device split → C++ frontend → NVVM IR (LLVM) → ptxas → SASS; what cudafe++, cicc, ptxas, fatbinary each do
• PTX: a virtual ISA — typed registers (infinite), explicit state spaces (.global, .shared, .param, .reg), predication, the %tid/%ctaid special registers
• SASS: the real ISA — register allocation, dual-issue, stall counts in control bits; why SASS changes every architecture but PTX doesn't
• JIT vs AOT: -arch/-code, compute_XX vs sm_XX, fatbins, the driver JIT cache (~/.nv/ComputeCache), forward compatibility — and the production incident pattern ("new GPU, 5-minute cold start")
• What the compiler does for you: loop unrolling, if-conversion to predication, FMA contraction, register allocation vs occupancy (-maxrregcount, the Phase 02 connection)
• The portability landscape a hardware-agnostic platform must navigate: LLVM NVPTX vs NVIDIA's NVVM, AMD's ROCm/LLVM path (GCN/RDNA ISA is public!), SPIR-V, Metal
• Triton in one lab-sized bite: block-level programming model, why it can fuse and auto-tune, where it fits between CUDA C++ and TVM/MLIR
• Compiler-as-moat: why kernel libraries + compiler maturity, not TFLOPs, decide which accelerators win (ties to Phase 01 Ch. 10, Phase 09)

Labs

Lab 01 — PTX/SASS Dissection (CUDA toolkit or Godbolt)

Lab 02 — Mini Compiler: Expression Language → PTX (Python)

Deliverables Checklist

• Fully annotated saxpy.ptx matching the answer key
• SASS findings: one predicated instruction, one FMA, one spill identified with flags that caused them
• Mini compiler passes all differential tests
• At least one optimization pass (folding or CSE) implemented and measured
• You can draw the nvcc pipeline and the JIT/AOT decision tree from memory

Interview Relevance

• "What does nvcc actually do? What's PTX vs SASS?"
• "Your service takes 5 minutes to start on new GPUs. Why, and what's the fix?" (JIT + fatbin story)
• "Why did Triton succeed where so many kernel DSLs failed?"
• "An accelerator startup says 'we run PyTorch.' What's actually behind that claim, and what do you audit?"
• "Walk me through compiling c[i] = a[i] * 2 + b[i] down to PTX by hand."

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
• Chapter 2: Compilers From Zero — the Five Classic Stages
• Chapter 3: The NVCC Pipeline, For Real
• Chapter 4: PTX — The Virtual ISA That Is CUDA's Moat
• Chapter 5: Reading PTX — A Guided Tour of saxpy
• Chapter 6: SASS — Where the Machine Stops Pretending
• Chapter 7: JIT vs AOT, Fatbins, and the Cold-Start Incident
• Chapter 8: What the Optimizer Does to Your Kernel
• Chapter 9: The Portability Landscape — NVPTX, ROCm, SPIR-V, Triton, MLIR
• Chapter 10: Compiler-as-Moat — The Business Reading
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

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)

🛸 Hitchhiker's Guide — Phase 03: GPU Compilers & Codegen

Read this if: you can write CUDA but nvcc is a magic box, "PTX" is a word you nod at, and you couldn't yet audit an accelerator vendor's "we support PyTorch" claim. Compressed field notes; derivations live in the WARMUP.

0. The 30-second mental model

CUDA C++ ──(cicc/LLVM)──► PTX (virtual ISA, stable forever)
                            └─(ptxas, AOT)──► SASS for sm_80, sm_90...   } fatbin
                            └─(driver JIT, load time)──► SASS for whatever ships next

PTX = the contract. SASS = the truth. Fatbin = the deployment policy. JIT = the cold-start you forgot to plan for. Everything else is detail.

1. The toolbox

2. PTX reading card

• %tid.x %ctaid.x %ntid.x = threadIdx / blockIdx / blockDim
• mad.lo.s32 %r4, %ctaid, %ntid, %tid = THE index formula
• .param → ld.param → cvta.to.global = kernel ABI prologue (your Lab 02 compiler emits this trio)
• setp.ge.s32 %p1, … + @%p1 bra = tail guard; @%p on a non-branch = if-conversion (no divergence machinery at all)
• fma.rn.f32 = contracted multiply-add; one rounding, not two — the reason for tolerance-based testing, and for -fmad=false in numerical disputes
• Infinite %f/%r/%rd registers — pressure is invisible here; spills only exist in SASS

3. Production policies you now own

• Build matrix: AOT SASS for every arch in the fleet + newest compute_XX PTX for the future. Cost: build minutes, container GB. Skipping an arch = JIT cold-start storm (or hard load failure if you also skipped PTX).
• JIT hygiene: CUDA_MODULE_LOADING=LAZY (default since 12.x), pre-warm compute cache in images, watch CUDA_CACHE_MAXSIZE on multi-model fleets.
• Numerical policy: FMA contraction + reduction order = CPU/GPU diffs in the last ulp; tests compare with tolerance, never == (Phase 02), and -fmad changes are reviewed like API changes.
• The vendor audit (the leadership move of this phase): "Top-20 ops for our workload: % of roofline at p50/p99, on our shapes and precisions. Which came from a kernel library, which from your compiler, which are unimplemented?" Watch them sweat on attention variants and odd head dims.

4. The portability map, one line each

• NVPTX (LLVM): how every non-NVIDIA frontend reaches NVIDIA hardware.
• ROCm/HIP: CUDA-shaped API, mechanical porting, public ISA — but no PTX equivalent: per-gfx code objects, no forward-compat promise. Plan rebuilds.
• SPIR-V: truly neutral, historically below peak for ML; the "standards" position.
• Triton: block-level Python DSL → MLIR → PTX/ROCm; ~90% of hand-tuned on LLM kernels; the pragmatic kernel layer for a hardware-agnostic platform.
• MLIR: the infrastructure bet; what you adopt so you don't invent an IR.

Three platform postures: wrap vendor libraries behind a HAL (fast, Phase 09), adopt Triton/MLIR for kernels (velocity), build a full compiler (moonshot). Blend of first two = the defensible default. Be ready to defend it to a board.

5. War stories to recognize

• "New A-series GPUs took 8 minutes to first token" — no sm_XX cubin, JIT storm. Fix: build matrix + cache pre-warm. (You will meet this one.)
• "Same model, different last digit on GPU vs CPU" — FMA contraction + parallel reduction order. Not a bug; a numerical policy.
• "We set maxrregcount=32 and got slower" — spills to local memory; the SASS shows STL in the inner loop. Occupancy isn't a score (Phase 01/02).
• "Vendor demo hit 80% of peak; our model hit 19%" — demo ran their three tuned kernels; your model hit the unimplemented long tail. The audit in §3 exists because of this meeting.

6. Exit bar

You can: narrate saxpy's PTX cold; find spills/FMA/predication in SASS; draw the fatbin/JIT decision tree and its incident; defend a build-matrix policy; argue the three portability postures; and your own little compiler emits PTX that a real driver loads. That's "technical credibility with engineering teams" at the compiler layer — on to Phase 04, where we descend below the runtime into the driver and the kernel (the OS one).

Lab 01 — PTX/SASS Dissection

Phase: 03 — GPU Compilers | Difficulty: ⭐⭐⭐☆☆ | Time: 3–4 hours Tools: CUDA toolkit (nvcc, cuobjdump) — or godbolt.org with zero install

Concept primer: ../WARMUP.md Ch. 3–8.

Run

nvcc -ptx -o saxpy.ptx kernels.cu                  # PTX for annotation
nvcc -cubin -arch=sm_80 -o kernels.cubin kernels.cu
cuobjdump --dump-sass kernels.cubin > kernels.sass # SASS for the hunt
nvcc -Xptxas=-v -c kernels.cu                      # registers/smem/spills

No toolkit? godbolt.org → language CUDA C++ → compiler "NVCC (latest)" → flags -O3 -arch=sm_80; the right pane shows PTX, and "Add new → Device" shows SASS. Everything below except the fatbin step works there.

0. The mission

Four tiny kernels (kernels.cu), four investigations:

1. saxpy — annotate every PTX line. Template in saxpy-annotation-template.md; answer key in saxpy-annotated.md. Do not peek early.
2. guarded — a small if body: find where the branch disappeared in SASS (if-conversion → @%p-predicated stores / SEL).
3. unrolled — #pragma unroll 8 loop: count the replicated bodies in SASS; note the register count change in -Xptxas=-v.
4. spiller — a kernel with a large per-thread array. Compile normally, then with -maxrregcount=16: find the LDL/STL local-memory spills that appear, and the "bytes spill stores/loads" line in ptxas output.

Then the fatbin dissection:

nvcc -gencode arch=compute_80,code=sm_80 \
     -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 -c kernels.cu -o fat.o
cuobjdump --list-elf fat.o     # two ELFs (sm_80, sm_90)...
cuobjdump --dump-ptx fat.o | head   # ...plus embedded PTX (compute_90)

Map each artifact to the load-time decision tree (WARMUP Ch. 7) and write, in two sentences, what happens on (a) an A100, (b) an H100, (c) whatever ships next.

1. The annotation rubric (what "done" means for saxpy)

Your annotation must identify: the .param ABI and ld.param prologue; the cvta.to.global state-space conversion and why it exists; the mad.lo.s32 index formula and which special registers feed it; the setp+@%p bra tail guard; the mul.wide.s32 byte-offset computation; the fma.rn.f32 contraction (and what -fmad=false would do to it). Six items — the same six the WARMUP Ch. 5 tour narrates.

2. Findings sheet (fill in)

3. Extensions

• Cross-compiler: same saxpy through clang -x cuda --cuda-device-only -S (LLVM NVPTX) — diff the PTX against nvcc's. What's identical (the ABI), what differs (instruction selection, virtual register naming)?
• Cross-vendor: on Godbolt, compile the HIP version with hipcc for gfx90a and skim the GCN ISA — note there's no PTX-like layer (WARMUP Ch. 9's asymmetry, seen with your own eyes).
• Tensor Core spotting: compile a 16×16 wmma example and find HMMA in SASS — the instruction the whole AI economy runs on.

4. What this lab proves about you

The compiler stack is no longer folklore: you've read the contract layer (PTX), the truth layer (SASS), and the deployment layer (fatbin) with your own eyes, and you can connect each to a production decision (numerical policy, occupancy tuning, build matrix).

Lab 02 — Mini Compiler: Expression Language → PTX

Phase: 03 — GPU Compilers | Difficulty: ⭐⭐⭐⭐☆ | Time: 6–10 hours Language: pure Python (stdlib) | Hardware: none — includes a PTX interpreter; optional real-GPU load at the end

Concept primer: ../WARMUP.md Ch. 2 (compiler stages), Ch. 4–5 (PTX), Ch. 8 (the optimizations you'll add).

Run

python solution.py          # compiles 5 kernels to PTX, runs differential tests
python solution.py --dump   # also prints every kernel's emitted PTX

0. The mission

Build the real thing, small: a compiler for elementwise GPU kernels —

kernel saxpy(a: f32, x: f32*, y: f32*) {
    y[i] = a * x[i] + y[i];
}

— that emits legal PTX (the same shape you annotated in Lab 01: .param ABI, cvta.to.global, mad.lo.s32 index, setp guard, fma.rn.f32), where i is implicitly the global thread index and a hidden n parameter guards the tail.

Because most readers don't have a GPU attached, the lab includes a PTX interpreter (PTXInterp) that executes the instruction subset per-thread with float32 rounding — so every compiled kernel is differentially tested against a plain-Python reference. Compiler people call this differential testing; it's how real codegen is validated (alongside conformance suites), and building the interpreter teaches the ISA semantics a second way.

1. The pipeline (read in this order)

Each stage is ~60 lines. The deep lesson: a compiler is a chain of small, testable transforms — after this lab, "we should build a compiler layer" stops being a scary sentence and becomes a costable one.

2. What the emitted PTX looks like

For c[i] = a[i] * 2.0 + b[i]; you'll get exactly the Lab 01 idioms:

.visible .entry scale2(
    .param .u64 scale2_param_0,      // a
    .param .u64 scale2_param_1,      // b
    .param .u64 scale2_param_2,      // c
    .param .u32 scale2_param_3       // n  (implicit)
){
    ...
    mad.lo.s32 %r4, %r1, %r2, %r3;   // i
    setp.ge.s32 %p1, %r4, %r5;
    @%p1 bra DONE;
    mul.wide.s32 %rd10, %r4, 4;      // byte offset, computed once
    ...
    ld.global.f32 %f1, [%rd11];
    mov.f32 %f2, 0f40000000;         // 2.0 as IEEE-754 hex — PTX float literal form
    ld.global.f32 %f3, [%rd12];
    fma.rn.f32 %f4, %f1, %f2, %f3;   // contraction: (a*2.0)+b -> one FMA
    st.global.f32 [%rd13], %f4;
DONE:
    ret;
}

Note 0f40000000: PTX spells float immediates as IEEE-754 bit patterns — one of a dozen small ABI/syntax facts you only learn by emitting the format for real.

3. Extension ladder (this is the lab)

1. Intrinsics: add min(x, y) / max(x, y) → min.f32/max.f32. Touches every stage once — the perfect first change.
2. Constant folding: a pass over the AST: 2.0 * 4.0 + x → 8.0 + x. Assert emitted instruction count drops.
3. CSE: (a[i]+b[i]) * (a[i]+b[i]) currently loads and adds twice. Hash subtrees, reuse registers. Measure again.
4. relu(a[i] * b[i] + c[i]) — fused elementwise kernel via max + folding: you have just built, in miniature, what fusion compilers (XLA, Inductor) do at scale.
5. Real hardware (optional, any NVIDIA GPU):

   from cuda import cuda  # pip install cuda-python
   # cuLinkCreate/cuModuleLoadDataEx accepts PTX text directly:
   err, mod = cuda.cuModuleLoadData(ptx_text.encode())
   

   Launch with cuLaunchKernel and compare against the interpreter. When the driver accepts your compiler's output, you've closed the loop this phase is about.
6. Second backend (Phase 09 foreshadowing): emit C source from the same AST (a for loop over i). Notice what's shared (front/middle end) and what isn't (ABI, index model) — that boundary is the HAL design problem.

4. Common pitfalls

1. Float literals as decimals — PTX wants 0f<bits> form for exactness; emitting 2.0 works in some toolchains and silently loses precision in others. The provided f32_hex() does it right.
2. Reusing the byte-offset register per operand — compute i*4 once; recomputing it per load is correct but bloats the kernel (and your CSE pass should catch it if you regress).
3. Forgetting the guard — the interpreter intentionally runs n not equal to a block multiple; an unguarded kernel writes out of bounds and the differential test catches the corruption. (Same bug as Phase 02 Lab 01 §1.)
4. f32 rounding in the oracle — the interpreter rounds every op to float32 (struct.pack round-trip); compare with relative tolerance, never == (FMA contraction changes the last ulp — WARMUP Ch. 8).

5. What this lab proves about you

"I built a compiler that emits PTX the CUDA driver loads, with differential tests" — that sentence, backed by this repo, is technical credibility with any compiler or kernels team you will ever lead. And the build-vs-adopt argument (WARMUP Ch. 9–10) is no longer abstract: you know exactly which parts were easy (the front end), which were fiddly (ABI details), and which you didn't even attempt (optimization at scale, multi-arch lowering) — which is precisely the cost model an engineering head needs.

Phase 04 — Linux Internals, GPU Drivers & Containers

Difficulty: ⭐⭐⭐⭐☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], Infrastructure/Platform Engineer, Driver-adjacent Runtime Engineer Hardware needed: a Linux machine or VM (Lima/Multipass/UTM on macOS, WSL2 on Windows, any cloud VM). No GPU required.

Why This Phase Exists

The JD requires "solid grounding in Linux internals, containers, GPU drivers." Between your CUDA program (Phase 02) and the silicon (Phase 01) sits a stack most engineers treat as superstition: libcuda.so, /dev/nvidia0, a kernel module, ioctls, mapped ring buffers. When a fleet GPU "falls off the bus," when a container can't see its GPU, when driver/CUDA version skew bricks a deployment at 2 a.m. — the engineers who can reason about this layer fix it; everyone else reboots and prays.

You will not write an NVIDIA driver. You will write a small, honest Linux character-device driver that exposes the same architectural pattern GPU drivers use — a /dev node, ioctl command submission, mmap-style buffer sharing — and you will build a container from scratch with namespaces and cgroups, then map exactly how a GPU passes through that boundary. After this, nvidia-container-toolkit and k8s device plugins (Phase 06) are obvious instead of magical.

Concepts

• Kernel space vs user space; syscalls; what a "driver" actually is (callbacks registered with the kernel)
• Character devices: major/minor numbers, /dev nodes, udev, file_operations
• ioctl: the universal "weird device command" escape hatch — and the GPU's actual control path
• The real NVIDIA stack: KMD (nvidia.ko) vs UMD (libcuda.so), command ring buffers, doorbells, why submission avoids syscalls in the hot path
• GSP firmware, the open-gpu-kernel-modules split, nouveau vs proprietary
• Driver/CUDA version skew: forward-compat matrix, the "driver too old" failure, fleet upgrade choreography
• DMA and IOMMU; pinned memory from the kernel's point of view (Phase 02 Ch. 5 connection)
• Namespaces (mount, PID, net, user) — what a container actually is
• cgroups v2: cpu, memory, io controllers — and why GPUs have no cgroup controller (the gap that creates the entire GPU-sharing industry, Phase 06)
• How a GPU enters a container: device nodes + library injection (nvidia-container-toolkit), CUDA_VISIBLE_DEVICES vs cgroup devices allow-lists
• /proc, /sys, strace, lsof, dmesg — the debugging toolkit

Labs

Lab 01 — A GPU-Shaped Character Device Driver (C, Linux kernel module)

Lab 02 — Container From Scratch + GPU Passthrough Anatomy (C + shell)

Deliverables Checklist

• fakegpu.ko loads; client passes all asserts; you can explain every file_operations callback
• Annotated strace: your client vs a real CUDA app (or the captured trace provided)
• Mini-runtime: isolation + OOM demos reproduced
• The "how a GPU enters a container" explainer written
• You can narrate the KMD/UMD split and the ring-buffer submission path from memory

Interview Relevance

• "What happens, at the OS level, when you call cudaLaunchKernel?"
• "Why doesn't CUDA_VISIBLE_DEVICES provide isolation? What does?"
• "A container sees the GPU but CUDA init fails. Debug it." (driver/library version skew, device node perms, the toolkit hook)
• "Why is there no GPU cgroup controller, and what do people do instead?"
• "Driver upgrade across a 10k-GPU fleet — walk me through the plan."

Warmup Guide — Linux Internals, GPU Drivers & Containers

Zero-to-expert primer for Phase 04. Assumes you can use a Linux shell and read C. No kernel programming background assumed — that's what we build. By the end the stack between cudaLaunchKernel and silicon is glass, not fog.

Table of Contents

• Chapter 1: Kernel Space, User Space, and What a Driver Is
• Chapter 2: Character Devices — /dev, major/minor, file_operations
• Chapter 3: ioctl — The GPU's Actual Control Path
• Chapter 4: The Real NVIDIA Stack — KMD, UMD, Ring Buffers, Doorbells
• Chapter 5: DMA, IOMMU, and Pinned Memory From the Kernel's Side
• Chapter 6: Version Skew — The Driver/CUDA Compatibility Matrix
• Chapter 7: Namespaces — What a Container Actually Is
• Chapter 8: cgroups v2 — Resources, and the GPU-Shaped Hole
• Chapter 9: How a GPU Enters a Container
• Chapter 10: The Debugging Toolkit — strace, dmesg, /proc, /sys
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: Kernel Space, User Space, and What a Driver Is

Zero background: the CPU runs in (at least) two privilege modes. Kernel mode can touch hardware, page tables, interrupts; user mode cannot — any attempt traps. Your process lives in user mode; the only doorway into kernel mode is a syscall (or an interrupt/fault). That hard boundary is why "the process crashed" doesn't take the machine down, and why everything hardware-related — including every GPU operation — ultimately funnels through the kernel at least once.

What a driver is, demystified: a driver is not a process or a daemon. It's a bundle of callbacks registered with the kernel ("when someone open()s this device, call my open; when they ioctl(), call my ioctl") plus interrupt handlers ("when the hardware raises IRQ N, call this"). A loadable kernel module (.ko) is just a relocatable object the kernel links into itself at runtime — insmod is closer to dlopen than to exec. Lab 01 makes this concrete: your fakegpu.ko is ~300 lines of callbacks and one struct file_operations.

The cost model that explains GPU driver design: a syscall costs ~100–300 ns plus cache pollution; an interrupt costs more. Drivers for high-rate devices (GPUs, NICs, NVMe) are therefore architected to get the kernel out of the hot path: set up shared memory once, then submit work from user space without syscalls (Ch. 4). Keep this in mind — it's the single design force behind everything in this phase.

Misconception: "the GPU driver runs on the GPU." The driver runs on the CPU, in the kernel; the GPU runs firmware (increasingly — GSP, Ch. 4) and your kernels. The driver's job is bookkeeping, memory management, scheduling, and shoveling command buffers.

Chapter 2: Character Devices — /dev, major/minor, file_operations

Unix's grand unification: devices are files. open("/dev/nvidia0") returns an fd; read/write/ioctl/mmap on that fd dispatch into driver callbacks.

The plumbing, end to end:

• A character device ("char dev") is the streams-of-bytes/commands flavor (vs block devices, which are the page-cached storage flavor). GPUs are char devices.
• Each device has a major number (selects the driver) and minor number (selects the instance — /dev/nvidia0 vs /dev/nvidia1). ls -l /dev/nvidia* shows them (major 195 for NVIDIA, historically).
• /dev nodes are created by udev at hotplug (or mknod by hand, or devtmpfs). A node is just a (type, major, minor) triple in a filesystem — this fact is load-bearing for containers (Ch. 9): injecting a GPU into a container is largely creating/allowing this node inside it.
• In the driver, you register a struct file_operations:

static const struct file_operations fakegpu_fops = {
    .owner          = THIS_MODULE,
    .open           = fakegpu_open,      // allocate per-fd state — "a context"
    .release        = fakegpu_release,   // free it
    .unlocked_ioctl = fakegpu_ioctl,     // the command path
    .mmap           = fakegpu_mmap,      // shared-memory path (extension)
};

The idea to carry away: per-open-file state. Each open() gets its own context structure (Lab 01 hangs a buffer registry and job queue off it). This is precisely what a CUDA context is at the driver level — and why two processes on one GPU don't see each other's allocations (and why a poisoned context dies alone, Phase 02 Ch. 4).

Chapter 3: ioctl — The GPU's Actual Control Path

read/write suit byte streams. Devices need commands: "allocate 2 GB," "map this buffer," "submit this command stream." The escape hatch is ioctl(fd, CMD, arg) — a generic "do device-specific thing," where CMD encodes direction/size/type and arg points to a command struct copied across the boundary.

The ABI discipline (Lab 01's core lesson):

struct fakegpu_alloc { __u64 size; __u32 handle; /* out */ };
#define FAKEGPU_IOC_ALLOC  _IOWR('F', 1, struct fakegpu_alloc)

// kernel side — NEVER dereference user pointers:
if (copy_from_user(&req, (void __user *)arg, sizeof(req))) return -EFAULT;
...validate everything...
if (copy_to_user((void __user *)arg, &req, sizeof(req))) return -EFAULT;

copy_from_user/copy_to_user are the checked crossings; validation is mandatory because user space is adversarial (a security boundary you'll revisit in Phase 11). Handles — small integers the driver maps to kernel-side objects — instead of raw pointers, for the same reason fds exist.

Proof that this is real: strace any CUDA program:

openat(AT_FDCWD, "/dev/nvidiactl", O_RDWR) = 3
ioctl(3, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2a, 0x20), 0x7ffd...) = 0
openat(AT_FDCWD, "/dev/nvidia0", O_RDWR)  = 4
ioctl(4, ...)                             = 0     // dozens at init
mmap(NULL, 0x200000, ..., fd=4, ...)      = 0x7f...   // ring buffers!

Initialization is a flurry of ioctls; then it goes quiet — because the steady state uses the mmap'd path (next chapter). Lab 01 step 3 has you place these two traces side by side.

Chapter 4: The Real NVIDIA Stack — KMD, UMD, Ring Buffers, Doorbells

The division of labor:

• KMD (kernel-mode driver) — nvidia.ko (+ nvidia-uvm.ko for unified memory, nvidia-modeset.ko for display): privileged setup. Memory management (GPU page tables!), context creation, channel/ring allocation, interrupt handling, recovery (Xid errors — Phase 10), power.
• UMD (user-mode driver) — libcuda.so (the driver API from Phase 02 Ch. 1 lives here): everything per-launch. Builds command buffers in memory shared with the GPU, manages your virtual address space, JIT-compiles PTX (Phase 03 Ch. 7!), implements the CUDA API surface.

The hot path — why launches take ~5 µs, not ~5 ms:

1. At context setup, the KMD maps a ring buffer (command queue) and a doorbell register into the process's address space (those mmaps in the strace).
2. To launch a kernel, libcuda.so writes the command into the ring (a memory write!) and writes the doorbell (another memory write, which the hardware sees) — zero syscalls.
3. The GPU's front end fetches commands, executes, and writes a fence (completion marker) back to shared memory; the UMD polls or sleeps on an interrupt for it. (cudaEventSynchronize = wait-on-fence.)

This user-space-submission pattern is universal in high-performance I/O — NVMe queues, io_uring, DPDK NICs, every modern GPU (AMD's is in the open amdgpu driver — same shapes, readable source). Your Lab 01 extension (mmap'd completion ring) implements the fence half.

Two modern wrinkles worth knowing: GSP firmware — recent NVIDIA drivers offload much KMD logic to a RISC-V management processor on the GPU, shrinking kernel-side code (and enabling the open-gpu-kernel-modules release: the kernel module is now MIT/GPL dual-licensed and on GitHub — readable!); the proprietary crown jewels moved to firmware + UMD. And nouveau — the reverse-engineered FOSS driver — historically lacked the firmware interfaces for reclocking, hence its performance gap; GSP is slowly changing that equation.

Chapter 5: DMA, IOMMU, and Pinned Memory From the Kernel's Side

Phase 02 Ch. 5 told you pinned memory makes transfers ~2× faster and truly async. Now the mechanism:

• DMA: the GPU's copy engines read/write host RAM directly, without the CPU touching the data. To do that safely, the physical pages must (a) stay resident — not be swapped or migrated — and (b) be known to the device.
• Pinning (cudaMallocHost → get_user_pages/page-locking in the kernel) guarantees (a). For pageable memory the driver must instead copy through an internal pre-pinned bounce buffer — there's your 2× and your fake-async.
• IOMMU provides (b)-with-safety: it's a page table for devices, translating device-visible addresses and blocking DMA outside allowed ranges. Cloud passthrough (vfio), confidential computing, and "why does iommu=off make it faster but scarier" all live here. (Security relevance returns in Phase 11: DMA from a malicious device is a classic attack class.)

Chapter 6: Version Skew — The Driver/CUDA Compatibility Matrix

The unglamorous chapter that pages people at 2 a.m.:

• One kernel module serves all processes; each container/app brings its own CUDA toolkit/runtime. The contract: driver ≥ what the toolkit needs. nvidia-smi's "CUDA Version" = the max the driver supports, not what's installed.
• Failure signatures to recognize on sight: CUDA driver version is insufficient for CUDA runtime version (driver too old); forward compatibility was attempted on non supported HW (compat package misuse); a container that worked on node A failing on node B (driver skew across the fleet); kernel upgrade → module won't load → DKMS didn't rebuild (or secure-boot signing failed — Phase 11 connection).
• Fleet upgrade choreography (the head-of-engineering deliverable): inventory driver/toolkit pairs → canary nodes → drain GPU workloads (Phase 06 scheduler integration) → upgrade KMD (reboot or careful unload) → health-check (Phase 10's checks) → progressive rollout. The forward-compatibility packages (newer libcuda on older driver branch, datacenter GPUs only) buy flexibility; pinning exact driver versions in node images buys reproducibility. Pick a policy and write it down — Lab 02's explainer and Phase 10's release engineering both cite this chapter.

Chapter 7: Namespaces — What a Container Actually Is

There is no container object in the kernel. "Container" = a process started with unshared namespaces + resource limits (Ch. 8) + (usually) reduced privileges. Namespaces virtualize visibility:

One syscall family creates them: clone(CLONE_NEWPID | CLONE_NEWNS | ...) / unshare(2). Lab 02's runtime is ~200 lines because that's genuinely all it takes: clone with flags, set hostname, pivot_root into a rootfs, mount fresh /proc, drop into a shell. Docker/Podman/containerd add image management, networking, and an API — the isolation itself is these kernel primitives.

Why this matters for a GPU platform: every scheduling, isolation, and security promise you make to customers (Phase 06, 11) bottoms out in which namespaces/cgroups/devices a workload got. You can't reason about "can tenant A see tenant B's model weights" without this chapter.

Chapter 8: cgroups v2 — Resources, and the GPU-Shaped Hole

Namespaces control what a process sees; cgroups control what it uses. v2 is a unified tree mounted at /sys/fs/cgroup; directories are groups; writes to control files set limits:

mkdir /sys/fs/cgroup/demo
echo 100M  > /sys/fs/cgroup/demo/memory.max
echo "50000 100000" > /sys/fs/cgroup/demo/cpu.max   # 0.5 CPU
echo $$    > /sys/fs/cgroup/demo/cgroup.procs       # join

Exceed memory.max → the OOM killer fires inside the group (Lab 02 step 2 demonstrates). The devices controller (an eBPF program in v2) allow-lists device-node access — the enforcement half of GPU injection (Ch. 9).

The GPU-shaped hole — say this sentence in interviews: there is no GPU cgroup controller. No gpu.max, no kernel-enforced fraction of SMs or HBM. Consequences:

• "GPU limits" in k8s are allocation (whole devices via the device plugin), not enforcement.
• Actual sharing requires GPU-side mechanisms — MPS (cooperative, weak isolation), MIG (hardware partition, strong), time-slicing (none) — Phase 06's entire subject matter, motivated by this one missing kernel feature.
• Vendors building "fractional GPU" products (and this JD's company likely among them) are filling exactly this hole with userspace interception or scheduler tricks — knowing why the hole exists (GPU contexts are opaque to the kernel; the KMD, not the scheduler, owns GPU time) is what lets you evaluate those products honestly.

Chapter 9: How a GPU Enters a Container

Now compose Chapters 2, 7, 8 — the full anatomy, which Lab 02 step 3 replicates manually:

A container needs three things to use a GPU:

1. The device nodes: /dev/nvidia0, /dev/nvidiactl, /dev/nvidia-uvm visible in its mount namespace (bind-mount or mknod), and allowed by the devices cgroup/eBPF program.
2. The user-mode driver libraries: libcuda.so matching the host kernel module (Ch. 6!). The container image cannot ship these — it would skew against the host KMD. So they're injected from the host at start.
3. No special kernel work — the KMD is shared; namespaces don't apply to it.

nvidia-container-toolkit is precisely this: an OCI hook (now CDI spec) that, at container creation, mounts the device nodes + host driver libs + nvidia-smi et al. into the container and writes the cgroup allow rule. CUDA_VISIBLE_DEVICES (or NVIDIA_VISIBLE_DEVICES at the hook level) then filters which minors get this treatment — note carefully: visibility, not security — a process that can mknod or escape the allow-list can reach other GPUs; the cgroup rule is the actual enforcement, and per-tenant isolation worthy of the name means dedicated devices or MIG (Phase 06) plus the Phase 11 story.

Debugging ladder for "container can't see GPU" (memorize; it's a weekly occurrence somewhere in any fleet): node has driver? (nvidia-smi on host) → toolkit hook configured? (runtime config) → nodes present in container? (ls /dev/nvidia*) → libs injected & matching? (ldconfig -p | grep cuda, compare versions) → cgroup allows? → only then blame the app.

Chapter 10: The Debugging Toolkit — strace, dmesg, /proc, /sys

The five tools that solve most layer-4 mysteries:

• strace -f -e trace=openat,ioctl,mmap — what is the process actually asking the kernel? (CUDA init failures become obvious: which open/ioctl returned what errno.)
• dmesg -T — the kernel's side of the story: module load errors, Xid errors (the GPU's error codes: Xid 79 = fell off the bus, Xid 48 = ECC DBE… Phase 10 builds alerting on these), OOM kills, IOMMU faults.
• /proc/PID/ — maps (see the mmap'd rings!), fd/ (which devices are open), status, cgroup.
• /sys — device topology (/sys/bus/pci/devices/...), cgroup tree, module params (/sys/module/nvidia/).
• lsof /dev/nvidia* — who's holding the GPU (the "why won't the driver unload" question).

Lab 01 step 3 and Lab 02's debugging exercises are drills on exactly these.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02 (Lab 02 injects Lab 01's device into a container — the punchline depends on both).

Lab 01 (fakegpu driver):

1. Read fakegpu.c top-down against Ch. 2–3 before building; predict what dmesg will print at insmod.
2. Build/load (any Linux VM; headers via apt install linux-headers-$(uname -r)). Run ./client; all asserts PASS.
3. strace ./client → annotate against the chapter-3 CUDA trace. Two columns, same shape — that's the lab's thesis made visible.
4. Extensions in order: quotas (validation discipline) → timer-based async completion (interrupt/fence flavor) → mmap'd completion ring (the real pattern; hardest).
5. Safety: it's a VM, snapshot first; a buggy module can panic the kernel — which is itself the lesson about why KMDs are small and UMDs are big.

Lab 02 (mini container runtime):

1. Run first, read after: sudo ./mininer run /bin/sh, poke around (ps, hostname, ls /), exit, then read mininer.c against Ch. 7.
2. cgroup demo: apply memory.max, run the included hog, watch the group OOM while the host stays healthy.
3. The synthesis step: inject /dev/fakegpu into the container (bind-mount + devices allow) and run Lab 01's client inside. When that passes, write the explainer — you have personally performed what nvidia-container-toolkit automates.

Success Criteria

• You can define "driver" as callbacks + interrupts, and narrate open→ioctl→copy_from_user→validate→copy_to_user (Ch. 1–3)
• You can explain why steady-state GPU submission uses zero syscalls (rings + doorbells + fences) and what the KMD vs UMD each own (Ch. 4)
• You can diagnose the four version-skew signatures and sketch the fleet upgrade choreography (Ch. 6)
• You can build a container from syscalls and state precisely what namespaces, cgroups, and capabilities each contribute (Ch. 7–8)
• You can recite the GPU-into-container anatomy and its debugging ladder, and you've performed it by hand (Ch. 9)
• "There is no GPU cgroup controller" — and you can derive Phase 06 from it

Interview Q&A

Q1: What happens at the OS level when a CUDA app starts and launches a kernel? A: Init: libcuda.so opens /dev/nvidiactl and /dev/nvidia0, issues a series of ioctls to create a context — the KMD allocates GPU page tables, a command ring, and maps the ring plus a doorbell into the process (visible as mmaps in strace). It also JIT-compiles any PTX lacking a matching SASS (Phase 03). Launch: the UMD writes the launch command into the mapped ring and taps the doorbell — two memory writes, no syscall — the GPU front end fetches, executes, and writes a fence; synchronization waits on that fence (poll or interrupt). That's why launches cost microseconds and why strace goes quiet after init.

Q2: Why doesn't CUDA_VISIBLE_DEVICES provide isolation? A: It's an environment variable read by the CUDA library in the untrusted process — it filters enumeration, nothing more; the process can unset it, and the device nodes are what they are. Enforcement lives lower: the devices cgroup/eBPF allow-list controls which nodes can be opened (what nvidia-container-toolkit configures), and even that only gates access to a whole device — within a shared GPU there's no kernel-enforced resource split because no GPU cgroup controller exists. Real isolation: dedicated GPU per tenant, or MIG hardware partitions, plus process-level (not context-level) multi-tenancy. The env var is UX; the cgroup is policy; MIG is enforcement.

Q3: A container sees /dev/nvidia0 but cudaInit fails. Debug. A: Ladder: (1) host sanity — nvidia-smi on the node (driver loaded? GPU healthy? check dmesg for Xids). (2) Library injection — ldconfig -p | grep libcuda inside the container; missing → toolkit hook not configured for this runtime; present → compare its version to the host KMD (nvidia-smi driver version) for skew, the classic "image ships its own libcuda" bug. (3) Permissions — devices cgroup allow rule, node major/minor match, non-root user in the right group. (4) The other nodes — nvidia-uvm missing breaks managed memory only, nvidiactl missing breaks everything. (5) strace the init and read which exact open/ioctl fails with which errno — it almost always names the culprit directly.

Q4: Why do GPU drivers split into a kernel module and libcuda.so, and why is the kernel part getting smaller? A: The split follows the privilege/performance line: things requiring kernel mode (page tables, interrupts, PCI config, cross-process arbitration) sit in the KMD; everything per-launch sits in user space for speed (ring submission) and shipability (a library updates with the toolkit; a kernel module update risks the host). The KMD is shrinking because GSP firmware on the GPU's own RISC-V core absorbed initialization and management logic — less GPL-adjacent kernel code (enabling NVIDIA's open kernel modules), faster context ops, and the proprietary logic now lives in firmware+UMD where it's both protected and updatable. Same pattern as NVMe/io_uring/DPDK: kernel for setup and safety, user space for the data path, firmware for vendor secrets.

Q5: Design the ioctl ABI for a simple accelerator. What discipline applies? A: Fixed-size, explicitly-padded structs (the ABI is forever — version field or reserved fields for growth); _IOWR macros encoding direction and size; copy_from_user, then validate every field against per-context state (handles not pointers, bounds on sizes, quota checks) because user space is adversarial; return handles, never kernel addresses; per-fd context so isolation falls out of the file model; errno discipline (-EINVAL vs -ENOMEM vs -EPERM mean different things to the caller's retry logic). And keep the hot path out of ioctl entirely — submission belongs in a mapped ring once setup is done (Lab 01's extension, and every real GPU's design).

Q6 (leadership): Plan a driver upgrade across a 10,000-GPU inference fleet. A: Inventory first: current driver/toolkit pairs per node pool, and every container image's CUDA runtime requirement — build the compatibility matrix before touching anything (driver must satisfy the max toolkit; forward-compat packages where datacenter GPUs allow). Canary: one node pool, full Phase 10 health suite (Xid monitoring, perf regression benchmarks — driver upgrades have shipped perf regressions). Choreography per node: cordon → drain via the scheduler (Phase 06; respect long-running jobs with checkpointing policy) → upgrade KMD + container toolkit together → reboot (cleanest; hot-unload is fragile with GPUs) → health-check → uncordon. Progressive rollout with automatic halt on Xid-rate or perf anomaly. Rollback path tested before rollout, image-pinned node pools so a rollback is a re-image, not an archaeology dig. And the boring decisive control: driver versions pinned in infrastructure-as-code, never apt upgrade'd by hand.

References

• Linux Device Drivers, 3rd ed. (Corbet, Rubini, Kroah-Hartman) — free online; Ch. 1–3, 6 map directly to Lab 01 — https://lwn.net/Kernel/LDD3/
• The kernel's own docs: The Linux Kernel documentation — driver-api, cgroup-v2 — https://docs.kernel.org
• NVIDIA open GPU kernel modules (read a real KMD!) — https://github.com/NVIDIA/open-gpu-kernel-modules
• NVIDIA Container Toolkit architecture + CDI spec — https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/
• NVIDIA CUDA Compatibility (the skew matrix) — https://docs.nvidia.com/deploy/cuda-compatibility/
• NVIDIA Xid errors reference — https://docs.nvidia.com/deploy/xid-errors/
• man 2 clone, man 7 namespaces, man 7 cgroups — genuinely excellent
• Liz Rice, "Containers from Scratch" (the famous live demo) — https://www.youtube.com/watch?v=8fi7uSYlOdc
• freedesktop.org DRM/amdgpu docs (the open-source GPU driver world) — https://docs.kernel.org/gpu/
• Cross-track: Phase 06 WARMUP (what fills the cgroup hole), security track (kernel/container hardening at depth)

🛸 Hitchhiker's Guide — Phase 04: Linux, Drivers & Containers

Read this if: CUDA works for you but /dev/nvidia0, nvidia.ko, and "the toolkit hook" are folklore. Field notes; derivations in the WARMUP.

0. The 30-second mental model

your app ─ libcudart ─ libcuda.so (UMD: rings, JIT, contexts)   [user space]
──────────────────────── ioctl/mmap on /dev/nvidia* ────────────[boundary]
nvidia.ko (KMD: page tables, channels, interrupts, recovery)    [kernel]
GSP firmware (management offload)  +  silicon                    [device]

Init = ioctl flurry + ring/doorbell mmaps. Steady state = zero syscalls: write command to mapped ring, tap doorbell, wait on fence. A container is a process with unshared namespaces + cgroup limits; a GPU enters it as device nodes + injected host libs. There is no GPU cgroup controller — that sentence generates Phase 06.

1. Commands that teach

strace -f -e openat,ioctl,mmap python -c "import torch; torch.zeros(1).cuda()"
ls -l /dev/nvidia*                      # major 195, minors = GPU index
dmesg -T | grep -i -e nvidia -e xid     # the kernel's side of the story
lsof /dev/nvidia0                       # who's holding the GPU
cat /proc/$(pgrep -f myapp)/maps | grep nvidia   # the mapped rings, visible
nvidia-smi -q | head -30                # driver version = max CUDA supported

2. The stack's division of labor

Rule that prevents half of all incidents: the container never ships libcuda.so; the host injects it. Images pinning their own = guaranteed eventual skew.

3. Container reality check

• Namespaces = visibility (mount/PID/net/UTS/user). cgroups = resources (cpu/memory/io + devices allow-list). Capabilities/seccomp = privilege. A "container" is just a process wearing all three.
• GPU "limits" in k8s are whole-device allocation, not enforcement. Fractions need MPS (cooperative), MIG (hardware, real), or time-slicing (none) — Phase 06.
• CUDA_VISIBLE_DEVICES is UX, the devices cgroup is policy, MIG is enforcement. Repeat before promising tenant isolation to a customer.

4. Failure signatures → causes (the 2 a.m. table)

5. War stories

• The image that shipped libcuda: passed CI (lucky match), died in prod region with older drivers. Policy since: images are scanned for driver libs; injection only.
• The quiet strace: engineer "proved" CUDA was hung because strace showed nothing — it was decoding correctly in the mmap'd ring the whole time. Steady state is silent by design.
• The OOM that wasn't: container kept dying at 100 MB despite "no limit" — inherited memory.max from a parent cgroup nobody remembered. cat /proc/self/cgroup first, always.
• The 40-minute driver upgrade that took a quarter: nobody had the image-pinned rollback path; one bad canary turned into fleet archaeology. WARMUP Q6 is the plan that avoids this.

6. Exit bar

You've written a driver (small, but real), built a container from syscalls, injected your own device into your own container, and can run the debugging ladder cold. The stack is glass. Phase 05 climbs back up one level: the runtime systems (allocators, stream schedulers) that live on top of this driver and under every inference engine.

Lab 01 — fakegpu: A GPU-Shaped Linux Character Device Driver

Phase: 04 — Linux & Drivers | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: C (kernel module + userspace client) | Hardware: Linux VM (snapshot it first!)

Concept primer: ../WARMUP.md Ch. 1–4. ⚠️ Kernel code runs with no safety net — use a VM (Lima/Multipass/UTM/WSL2/cloud), snapshot before insmod. A panic here is a lesson, not a disaster.

Run

sudo apt install build-essential linux-headers-$(uname -r)   # Debian/Ubuntu
make                       # builds fakegpu.ko (kernel) and client (user)
sudo insmod fakegpu.ko
dmesg | tail -3            # "fakegpu: loaded, /dev/fakegpu ready"
sudo ./client              # exercises the full ioctl ABI; prints PASS/FAIL
strace -e ioctl,openat ./client 2>&1 | head -30   # compare to a real CUDA trace
sudo rmmod fakegpu

0. The mission

Implement the architectural skeleton every GPU driver shares, in ~300 honest lines:

The client is the test suite: it asserts context isolation (two fds can't see each other's handles), quota enforcement (over-alloc returns ENOSPC), seqno ordering, and input validation (bad handle → EINVAL, oversized request → rejected — the adversarial-userspace discipline from WARMUP Ch. 3).

1. Reading order (fakegpu.c)

1. The ABI header (fakegpu_ioctl.h) — shared by kernel and client; note _IOWR encodings and fixed-size structs. ABI files are forever; this one is written like it.
2. fakegpu_open/release — per-fd context allocation. This is the pattern: everything else hangs off file->private_data.
3. fakegpu_ioctl — the dispatcher: copy_from_user → validate → act → copy_to_user. Count the validation lines vs the "work" lines (the ratio is the lesson).
4. The job model — SUBMIT enqueues and immediately completes jobs (incrementing completed_seq). Real hardware completes asynchronously; the extension adds a timer_list to simulate that.

2. The strace comparison (step 3 of the lab)

Your client:

openat(AT_FDCWD, "/dev/fakegpu", O_RDWR)       = 3
ioctl(3, FAKEGPU_INFO, ...)                    = 0
ioctl(3, FAKEGPU_ALLOC, {size=1048576, ...})   = 0
ioctl(3, FAKEGPU_SUBMIT, ...)                  = 0
ioctl(3, FAKEGPU_WAIT, ...)                    = 0

A real CUDA init (from WARMUP Ch. 3): same shape — open device, info/version ioctls, allocation ioctls, then mmap calls you don't have (the ring/doorbell — because real drivers move submission out of ioctl; your extension 3 closes that gap). Write the side-by-side into TRACE-NOTES.md — that artifact is your "technical credibility" exhibit for this phase.

3. Extensions

1. Quota tightening: per-context and global device quota; correct errno for each (ENOSPC vs ENOMEM). Add a /sys/module/fakegpu/ or sysfs stats attribute (your Phase 10 metrics on-ramp).
2. Async completion: a timer_list completes jobs 10 ms after submit; WAIT blocks on a wait_queue_head_t until completed_seq >= target. You've just implemented a fence.
3. The mmap ring (hard, the real pattern): fakegpu_mmap exposes a page containing completed_seq; client polls it with zero syscalls in steady state. Measure waits/sec: ioctl-WAIT vs mmap-poll (~100× — the entire reason GPU submission works the way it does).
4. Adversarial client: write evil.c — bad handles, huge sizes, ioctls on a closed-and-reopened fd, two threads racing SUBMIT. Every case must fail cleanly. (Phase 11 will call this "attack surface review"; here it's just good driver hygiene.)

4. Common pitfalls

1. Dereferencing the user pointer instead of copy_from_user — works in your VM (same address space mapping luck), then oopses elsewhere. The compiler __user annotation + sparse exists for this.
2. Forgetting mutex_lock around context state — two threads on one fd corrupt the buffer list; the racing-SUBMIT test in extension 4 catches it.
3. Leaking contexts on release — rmmod then dmesg shows the leak warning the module prints; the kernel's own kmemleak would find it in real review.
4. Building against the wrong headers (uname -r mismatch after a kernel update) — insmod: invalid module format. This is the DKMS lesson from WARMUP Ch. 6, self-inflicted once for learning.

5. What this lab proves about you

You can read (and write) the layer most GPU engineers only gesture at: when a fleet incident says "Xid 31, context poisoned, ioctl returned EIO," every word of that sentence now maps to code you've written. And nvidia's open-gpu-kernel-modules repo is now readable to you — start with kernel-open/nvidia/nv.c and recognize the shapes.

Lab 02 — Container From Scratch + GPU Passthrough Anatomy

Phase: 04 — Linux & Containers | Difficulty: ⭐⭐⭐⭐☆ | Time: 4–6 hours Language: C + shell | Hardware: Linux VM with root (same VM as Lab 01)

Concept primer: ../WARMUP.md Ch. 7–9.

Run

make                         # builds mininer + hog
sudo ./mkrootfs.sh           # builds ./rootfs from busybox (one download)
sudo ./mininer run /bin/sh   # you are now PID 1 in an isolated container
# inside: ps  /  hostname  /  ls /   — observe isolation
exit

sudo ./mininer run-limited 100 /bin/sh    # same, with memory.max=100M
# inside: /hog 200      — allocates 200 MB -> OOM-killed at ~100 MB

0. The mission

Build the smallest honest container runtime — then put a "GPU" inside it by hand, so that nvidia-container-toolkit becomes a thing you could have written.

mininer.c (~200 lines) does, in order:

1. clone(CLONE_NEWUTS | CLONE_NEWPID | CLONE_NEWNS | SIGCHLD) — new hostname, PID, and mount namespaces for the child.
2. In the child: sethostname("minibox"); make mounts private (MS_REC|MS_PRIVATE — so our mounts don't leak to the host); chroot("./rootfs") + chdir("/") (honest note: real runtimes use pivot_root, which unmounts the old root rather than hiding it — the extension asks you to upgrade); mount a fresh /proc (so ps sees only the namespace).
3. (run-limited) create /sys/fs/cgroup/mininer, write memory.max, add the child PID to cgroup.procs.
4. execv the requested command.

That's a container: namespaces (visibility) + cgroups (resources) + a rootfs (filesystem). Docker adds images, networking, and an API — not different physics.

1. The observation checklist (do these inside the container)

On the host, in another terminal: ps aux | grep sh still shows the container shell (as a normal PID) — namespaces virtualize the container's view, not the host's. Also cat /sys/fs/cgroup/mininer/memory.current while the hog runs.

2. The GPU passthrough synthesis (the point of the whole phase)

With Lab 01's module loaded on the host:

# 1. Device node into the container's rootfs (what the toolkit's hook does):
sudo mknod ./rootfs/dev/fakegpu c $(grep fakegpu /proc/misc | awk '{print $1}') \
     $(grep fakegpu /proc/misc | awk '{print $1}' > /dev/null; \
       ls -l /dev/fakegpu | awk '{print $6}' | tr -d ',')
#    (or simply: sudo cp -a /dev/fakegpu ./rootfs/dev/  — preserves maj:min)
sudo chmod 666 ./rootfs/dev/fakegpu

# 2. The "driver library injection" analog: copy the client in
sudo cp ../lab-01-fakegpu-driver/client ./rootfs/bin/gpuclient

# 3. Run it inside:
sudo ./mininer run /bin/gpuclient        # ALL CHECKS PASSED — inside a container

Now write GPU-IN-CONTAINER.md (one page): map what you just did to the three requirements of WARMUP Ch. 9 (device node + cgroup allowance + injected userspace libs; note that our mininer doesn't set up a device cgroup — add it as the extension and show /dev/fakegpu works while a non-allow-listed node fails). Close with the debugging ladder. This document is the deliverable — it's the team-training artifact the phase README promises.

3. Extensions

1. pivot_root upgrade: replace chroot with pivot_root + umount2(old, MNT_DETACH); explain in a comment why chroot alone is escapable (classic chroot jailbreak: keep an fd to the outside, fchdir).
2. Devices cgroup: attach a cgroup-v2 eBPF device filter (or use v1 devices.allow if available) permitting only /dev/fakegpu and /dev/null; demonstrate mknod + open of anything else failing.
3. Network namespace: CLONE_NEWNET + a veth pair (ip link add ... netns) — ping from container to host.
4. User namespace: CLONE_NEWUSER with uid_map so the container's root is unprivileged outside — then re-run the GPU test and document what breaks (device access vs UID mapping is exactly the rootless-container GPU pain point in real deployments).
5. Compare with runc: strace -f -e unshare,clone,mount,pivot_root runc run on a real container; diff against strace -f ./mininer run.

4. Common pitfalls

1. Forgetting MS_PRIVATE — your container's mounts propagate to the host (systemd makes / shared by default); the host's /proc gets shadowed. Symptom: weird host ps after exit. The MS_REC|MS_PRIVATE remount is non-negotiable.
2. Mounting /proc before the PID namespace exists — ps shows host processes; the fresh /proc must be mounted by the child, after clone.
3. cgroup v1 vs v2 confusion — this lab assumes v2 (/sys/fs/cgroup with memory.max). Check mount | grep cgroup2; hybrid hosts need the v2 path.
4. Running on your laptop instead of the VM — nothing here should damage a host, but you're root, juggling mounts; the VM habit from Lab 01 stands.

5. What this lab proves about you

You can state — and demonstrate, with code you wrote — exactly what a container is, what GPU injection consists of, and where the isolation boundaries actually sit. When a customer asks "how is tenant A isolated from tenant B on shared nodes," your answer comes from mechanism, not marketing. Phase 06 builds the scheduling layer on top of this; Phase 11 hardens it.

Phase 05 — Runtime Systems: Allocators, Streams, and Async Execution

Difficulty: ⭐⭐⭐⭐☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], Runtime Systems Engineer, ML Framework Engineer Hardware needed: none — both labs are CPU-only models of GPU runtime components (Rust + C++)

Why This Phase Exists

The JD asks for "runtime systems programming" by name. Between the driver (Phase 04) and the serving engine (Phase 07) lives the runtime: the layer that decides where memory comes from and when work runs. Every ML framework ships one (PyTorch's caching allocator, CUDA's stream machinery, TensorRT's execution contexts), and every inference platform's worst incidents — fragmentation OOMs at 60% utilization, deadlocked streams, mysterious sync stalls — happen here.

You will build the two canonical components yourself: a caching memory allocator in Rust (the PyTorch design: size-binned free lists, block splitting, stream-aware reuse) and a stream/event scheduler in C++ (dependency DAGs, ordered queues, fences). Both are faithful CPU-side models of the GPU originals, with tests that demonstrate the failure modes you'll spend your career preventing.

Concepts

• Why cudaMalloc/cudaFree are slow (device synchronization!) and fragmenting — the motivation for every framework allocator
• Allocator zoology: bump, buddy, slab, size-binned free lists; what malloc/jemalloc do vs what GPU pools need
• The PyTorch caching allocator design: bins, block splitting, rounding, empty_cache, fragmentation vs reserved-vs-allocated metrics, expandable segments
• Fragmentation: internal vs external; why long-running serving processes OOM at 60% "utilization"; defragmentation options (none, really — and what that implies for design)
• Stream semantics: ordered queues, default-stream legacy behavior, events as cross-stream edges (cudaStreamWaitEvent)
• The dependency DAG: how frameworks order kernels, copies, and frees; stream-ordered allocation (cudaMallocAsync) and why "free" is an event in a stream, not a moment in time
• Use-after-free across streams — the classic GPU runtime bug class and how stream-aware allocators prevent it
• CUDA Graphs: amortizing the ~5–10 µs launch cost; capture/replay semantics and their restrictions (why serving engines capture decode steps)
• Pinned-memory pools and H2D/D2H staging pipelines (Phase 02 Ch. 9, productionized)

Labs

Lab 01 — Caching GPU Allocator (Rust)

Lab 02 — Stream & Event Scheduler (C++)

Deliverables Checklist

• Allocator: all tests pass; fragdemo table reproduced and explained
• You can sketch the bin/split/cache design and the reserved-vs-allocated distinction from memory
• Scheduler: all scenarios pass; you can explain each failure mode it demonstrates
• One extension implemented per lab
• One page: "what our serving platform requires from its runtime layer" (feeds Phase 12 capstone)

Interview Relevance

• "Why does PyTorch have its own allocator? Walk me through its design."
• "nvidia-smi says 78 GB used but the framework reports 51 GB. Explain."
• "Your serving process OOMs after 3 days at constant load. Hypotheses?" (fragmentation, pinned leak, cache growth)
• "Design stream-ordered deallocation. What goes wrong without it?"
• "When do CUDA Graphs help, and what can't they capture?"

Warmup Guide — Runtime Systems: Allocators, Streams, Async Execution

Zero-to-expert primer for Phase 05. Assumes Phases 02 and 04 (the CUDA memory API and the driver beneath it). No allocator or concurrency background assumed. By the end you can design the runtime layer of an inference platform — and debug the ones you didn't design.

Table of Contents

• Chapter 1: The Runtime Layer — What Sits Between Driver and Engine
• Chapter 2: Why cudaMalloc Is Slow and Fragmenting
• Chapter 3: Allocator Zoology — Bump, Buddy, Slab, Bins
• Chapter 4: The Caching Allocator — PyTorch's Design, Dissected
• Chapter 5: Fragmentation — The Silent Fleet Killer
• Chapter 6: Streams and Events — The Ordering Algebra
• Chapter 7: Stream-Ordered Lifetime — Why "Free" Is an Event
• Chapter 8: CUDA Graphs — Amortizing Launch Overhead
• Chapter 9: Pinned Pools and Staging Pipelines
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: The Runtime Layer — What Sits Between Driver and Engine

The stack so far: silicon (Phase 01) ← kernels (02) ← compiler (03) ← driver (04). This phase is the layer that every framework and serving engine builds on top of the driver and underneath the model code:

model / engine code (vLLM, PyTorch, TensorRT)        Phase 07
─────────────────────────────────────────────
RUNTIME: memory pools, stream scheduling,            ← THIS PHASE
         graphs, staging pipelines
─────────────────────────────────────────────
CUDA runtime/driver API                              Phase 02/04

Why it exists at all: the driver's primitives (cudaMalloc, raw streams) have costs and semantics that are unacceptable to call in a hot loop. The runtime layer turns them into amortized, policy-bearing services. The two services that matter most — and that you'll build — are the allocator (Lab 01) and the execution scheduler (Lab 02).

Production significance: when this layer is wrong you get the worst class of incidents — not crashes, but slow decay: fragmentation OOM on day 3, throughput sag from false serialization, once-a-week corruption from a cross-stream lifetime bug. Nothing in a stack trace points here; you find these only if you know the layer exists.

Chapter 2: Why cudaMalloc Is Slow and Fragmenting

Three facts about the raw allocation API, each with a design consequence:

1. cudaFree synchronizes the device (historically; and cudaMalloc can too). The driver must be sure no in-flight kernel still uses the region — lacking fine-grained tracking, it waits for everything. One stray cudaFree per request = your entire pipeline gains a global barrier. → Consequence: never call it in steady state; allocate from a pool.
2. It's a driver round-trip (ioctl, page-table updates — Phase 04 Ch. 4): tens of µs vs ~100 ns for a pool hit. At thousands of tensor allocations per step, that's the difference between negligible and dominant. → Consequence: caching — keep freed blocks and reuse them.
3. There is no paging and no compaction. Host malloc lives on virtual memory: fragmented physical pages don't matter because the VM layer remaps. GPU allocations are (effectively) physically contiguous reservations in a fixed-size HBM; nothing remaps behind your back, and you cannot move a live block (kernels hold raw pointers). → Consequence: fragmentation is permanent damage within a process lifetime; the allocator's only weapons are placement policy and shape discipline (Ch. 5).

(cudaMallocAsync, CUDA 11.2+, moved a pool into the driver — same design, different owner; the VMM API (cuMemCreate + address reservation) added remapping tricks that "expandable segments" exploit — Ch. 4.)

Chapter 3: Allocator Zoology — Bump, Buddy, Slab, Bins

The four designs every systems engineer should hold (Lab 01 implements the fourth):

• Bump allocator: a pointer that only moves forward; free is a no-op (or whole-arena reset). O(1), zero fragmentation within an arena, but only fits phase-structured lifetimes. GPU relevance: activation memory in inference is exactly phase-structured — per-step arenas are a legitimate design (and what "static planning" compilers like TensorRT effectively do).
• Buddy allocator: powers-of-two blocks, split on demand, coalesce with your "buddy" on free. O(log n), bounded external fragmentation, up to ~50% internal waste. The classic kernel-page-allocator design.
• Slab allocator: per-size-class object caches. Brilliant for fixed-size objects — GPU relevance: KV-cache blocks (Phase 07) are fixed-size by design partly so a slab-style pool can manage them trivially.
• Size-binned free lists with splitting (the caching allocator): free blocks indexed by size class; allocation finds the smallest fitting block, splits the remainder back into the lists. This is malloc's general shape and PyTorch's choice — flexible sizes, good reuse, but external fragmentation is back (Ch. 5).

The meta-lesson: allocator design = exploiting what you know about lifetimes and sizes. The more structure you can impose on the workload (fixed block sizes, phase lifetimes, stream affinity), the better the allocator can be. This is why Phase 07's PagedAttention is best understood as an allocator redesign, not an attention algorithm.

Chapter 4: The Caching Allocator — PyTorch's Design, Dissected

The reference design (Lab 01 is a faithful Rust model):

• Two regimes by size: small (<1 MB) and large allocations get separate pools/bins (mixing tiny and huge requests is a fragmentation accelerant).
• Rounding: sizes round up to multiples (512 B; larger granularity for big blocks) — bounded internal waste purchases massive reuse improvement, because "1000 different sizes" collapses to "30 size classes."
• Allocation path: exact-ish bin hit → reuse (fast path, ~100 ns). Miss → take a larger cached block and split it (remainder returns to bins). Still nothing → cudaMalloc. OOM → free all cached-unsplit blocks, retry (the "sync-and-retry" you see as a mysterious stall right before a PyTorch OOM error).
• Free path: blocks go back to the bins — the driver never sees the free (this is why nvidia-smi shows "used" memory that the framework considers free: reserved vs allocated, the #1 memory-confusion in every ML org — Q2 below).
• Stream tagging: every block remembers the stream it was used on; reuse on a different stream requires an event dependency (or is simply not done) — preventing the cross-stream use-after-free (Ch. 7).
• expandable_segments (modern): uses the CUDA VMM API to grow a virtual reservation with physical pages on demand — restoring a little of the "remapping" power CPUs get for free, specifically to fight fragmentation.

Misconception to kill: "the allocator caches because cudaMalloc is slow." Half right — the deeper reason is cudaFree's synchronization and the fragmentation ratchet. Even with a free cudaMalloc you'd still want the pool.

Chapter 5: Fragmentation — The Silent Fleet Killer

Internal fragmentation: waste inside blocks (rounding). Bounded, boring, fine.

External fragmentation: free memory exists but in pieces too small for the request. The serving-fleet scenario (Lab 01's fragdemo reproduces it):

1. Day 0: requests of mixed sizes allocate/free happily; memory is a clean checkerboard.
2. Day 2: long-lived allocations (a resident model, a stuck request's KV) sit between short-lived ones. Frees leave holes of assorted sizes.
3. Day 3: a 2 GB request arrives. Free memory totals 9 GB. No hole exceeds 1.1 GB. OOM at 60% utilization. Restart fixes it (fresh address space) — which is why "we restart the pods nightly" is a real, embarrassing, industry-wide mitigation.

Because compaction is impossible (live raw pointers), the defenses are all preventive:

• Shape discipline: bucket request sizes (pad to standard lengths) — fewer distinct sizes = checkerboard stays clean. (Same trick as batching same-length sequences — Phase 01 Ch. 4's divergence logic, reborn.)
• Pool segregation: small/large pools; per-purpose pools (weights vs activations vs KV); per-stream pools.
• Fixed-size paging for the dominant consumer: if one consumer (KV cache) dominates and grows unpredictably, give it fixed-size blocks + an indirection table = external fragmentation eliminated for that consumer by design. That sentence is PagedAttention (Phase 07), derived from first principles.
• Metrics that see it coming (Phase 10): track reserved, allocated, and largest-free-block — the third is the early-warning signal nobody exports by default.

Chapter 6: Streams and Events — The Ordering Algebra

The execution model that Lab 02 implements, reduced to algebra:

• A stream is a FIFO: operations issued to the same stream execute in issue order. Within a stream you never need explicit sync.
• Different streams have no order — unless you create one.
• An event is a marker recorded into a stream; streamWaitEvent(s2, e) makes s2 pause until e fires. Events are the edges of the DAG; streams are convenient paths through it.
• cudaStreamSynchronize/cudaDeviceSynchronize join the CPU to the DAG — blunt instruments; in engine code, almost every device-wide sync is a bug or a missed event edge.
• The legacy default stream synchronizes with all other streams (a global serializer hiding in every naive codebase — Lab 02 scenario 2 measures the damage; the fix is per-thread default streams or explicit streams everywhere, Phase 02 Ch. 9).

Mental model for review: draw the intended DAG, then check the code creates exactly those edges — extra edges = lost overlap (performance bug); missing edges = race (correctness bug). Most stream code review reduces to this.

Chapter 7: Stream-Ordered Lifetime — Why "Free" Is an Event

The bug class that motivates stream-aware allocators (Lab 02 scenario 3):

stream A: kernel K reads buffer B          (still running...)
host:     free(B)                          (CPU is ahead of the GPU!)
allocator: hands B's memory to stream C    -> C writes over K's input. Corruption.

The CPU runs ahead of the GPU (async launches, Phase 02 Ch. 1) — so a host-side free() happens-before the GPU work that uses the buffer. The fix is to make deallocation an operation in the stream: the block returns to the pool only when the stream reaches the free point (an event). Reuse on another stream inserts a wait-event edge first. cudaMallocAsync/ cudaFreeAsync are exactly this, productized; PyTorch's stream tagging (Ch. 4) is the same idea retrofitted.

The leadership phrasing: lifetime on a GPU is defined by stream position, not by wall-clock or host program order. Any API in your platform that forgets this will eventually corrupt a customer's tokens — and it will be unreproducible, because it needs a scheduling coincidence. (This is the bug class behind a good fraction of "model produced garbage once" tickets.)

Chapter 8: CUDA Graphs — Amortizing Launch Overhead

The numbers (Phase 02): each kernel launch costs ~5–10 µs of CPU-side work. An LLM decode step is hundreds of short kernels; at high batch, GPU work per kernel can be tens of µs — the CPU becomes the bottleneck issuing work, and the GPU starves between kernels.

CUDA Graphs: record the whole sequence of launches (+ their dependency edges) once — capture — then replay the graph with one API call. Launch cost amortizes from n × 8 µs to ~1 × 10 µs + tiny per-node cost. Serving engines (vLLM et al.) capture the decode step per batch-size bucket; it's worth 10–30% at small batch and is the standard answer to "why is my decode CPU-bound?"

Restrictions (the part people learn the hard way): the graph is static — shapes, pointers, and control flow are baked at capture; capture forbids synchronous APIs (cudaMalloc!) inside the captured region (stream-ordered allocation exists partly for graphs); dynamic batch = capture per bucket or use graph update APIs. Lab 02's capture/replay extension makes all of this tactile in 100 lines.

Chapter 9: Pinned Pools and Staging Pipelines

The host↔device data path, productized (composing Phase 02 Ch. 5/9 with this phase's pooling):

• Pinned memory is mandatory for true async copies but expensive to allocate and steals unpageable RAM → pool it like device memory (a pinned staging-buffer pool is ~50 lines on top of Lab 01's design).
• The standard pipeline: chunk → H2D copy on copy-stream → kernel on compute stream (event edge) → D2H on copy-stream (event edge) — double-buffered so chunk N+1's copy overlaps chunk N's compute. GPUs have separate copy engines precisely so this overlap is real.
• Serving-specific instances you'll meet in Phase 07: weight loading at startup (pinned + multi-stream = minutes to seconds), KV-cache swap to host (the 50× bandwidth cliff makes the policy — what to swap when — the hard part, not the mechanism).

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02 (the scheduler's stream-ordered-free scenario reuses allocator concepts).

Lab 01 (allocator, Rust):

1. cargo test green first; then read src/lib.rs with Ch. 3–4 open. The types are small: Block, BinKey, CachingAllocator.
2. Trace one allocation by hand: 1.2 MB request → which bin, which split, what's left where. Write the trace in comments; the test test_split_accounting checks your understanding.
3. Run cargo run --release --bin fragdemo: the naive-vs-caching utilization table is the artifact. Explain the gap using Ch. 5's checkerboard story.
4. Extensions in order: coalescing (hard correctness) → empty_cache → trace replayer (turns the allocator into a measurable system — feed it your own traces).

Lab 02 (scheduler, C++):

1. make && ./scheduler_test; then read scheduler.hpp with Ch. 6–7 open — Stream is a worker+FIFO, Event is a counter+condvar, that's it.
2. Scenario 1 (overlap): predict the speedup from the DAG before reading the assert. Scenario 2 (legacy default stream): observe the global serializer. Scenario 3 (lifetime): the unsafe path corrupts; the stream-ordered path can't — connect each line to Ch. 7's diagram.
3. Extension: capture/replay. Measure issue-cost per op before/after — your own numbers for the Ch. 8 argument.

Success Criteria

• You can give the three reasons raw cudaMalloc/Free are unusable in steady state, with the design consequence of each (Ch. 2)
• You can sketch the caching allocator (bins, split, stream tags, sync-and-retry) and explain reserved vs allocated (Ch. 4)
• You can tell the day-3 fragmentation story with numbers and name the four preventive defenses (Ch. 5)
• You can review stream code by drawing its DAG and checking edges (Ch. 6)
• You can explain why free must be stream-ordered and what bug class ignores it (Ch. 7)
• You know what CUDA Graphs buy, their restrictions, and when serving engines use them (Ch. 8)
• Both labs pass; one extension each implemented

Interview Q&A

Q1: Why does PyTorch implement its own allocator instead of calling cudaMalloc? A: Three driver-level facts force it: cudaFree device-synchronizes (a global barrier per free), cudaMalloc is a driver round-trip (tens of µs vs ~100 ns pool hit), and HBM has no paging/compaction so fragmentation is permanent. The caching allocator answers all three: it never returns memory to the driver in steady state (frees go to size-binned lists), reuses blocks in ~O(1), rounds sizes to bound fragmentation, splits larger cached blocks for fit, tags blocks by stream to keep reuse safe under async execution, and on apparent-OOM frees its caches and retries. cudaMallocAsync later moved the same design into the driver as stream-ordered pools.

Q2: nvidia-smi shows 78 GB used; the framework says 51 GB. Who's lying? A: Nobody — they measure different lines of the same ledger. nvidia-smi sees what the driver handed the process: reserved memory, which includes everything the caching allocator is holding in its free lists. The framework's 51 GB is allocated — blocks currently handed to live tensors. The 27 GB gap is cache (plus context/fragmentation overhead). It's reclaimable via empty_cache() (at a performance cost) and is not a leak. The follow-up metric a platform should actually watch is largest free block — the fragmentation early-warning — and allocated-vs-reserved divergence trend, which distinguishes cache growth from a real leak.

Q3: A serving process OOMs after 3 days at steady load. Hypotheses, in order? A: (1) External fragmentation — mixed request sizes + long-lived allocations checkerboarding HBM until no hole fits a big request; check largest-free-block trend and whether OOM happens at <70% allocated; mitigate with size bucketing, pool segregation, paged KV (or, honestly, scheduled restarts while you fix it). (2) Genuine slow leak — allocated (not just reserved) trending up: a cache without eviction, per-request tensors retained by a logger, growing CUDA Graph pools per shape bucket. (3) Pinned host memory exhaustion masquerading as device OOM in mixed stacks. (4) KV-cache admission policy allowing pathological long-context requests to accrete. The diagnostic discipline: reserved vs allocated vs largest-block, per pool, on a dashboard before day 3 (Phase 10).

Q4: Two streams, one buffer: kernel on stream A reads it, host frees it, stream C gets it and writes. How do you make this class of bug impossible? A: Make deallocation stream-ordered: free() doesn't return the block to the pool — it records an event at the free point in A; the block becomes reusable only when that event completes. Reuse on a different stream additionally inserts streamWaitEvent so C's first use happens-after A's last. This is exactly cudaFreeAsync semantics / PyTorch's stream tagging. The general principle: on an async device, lifetime is a position in the dependency DAG, not a moment in host time — any API that takes raw "free now" must be wrapped before it touches engine code. Detection for the residue: compute-sanitizer in CI and a debug allocator mode that poisons freed blocks.

Q5: When do CUDA Graphs help, and what are the sharp edges? A: They help when launch overhead is material: many short kernels per iteration with a stable structure — the LLM decode step is the canonical case (hundreds of kernels, µs-scale each); capture per batch-bucket and replay, typically 10–30% at small batch and a fix for CPU-bound issuance. Sharp edges: the graph freezes shapes, addresses, and topology (dynamic control flow can't be captured); no synchronous calls (cudaMalloc) inside capture — you need stream-ordered allocation; per-shape graph pools consume memory (a real OOM contributor at high bucket counts); and debugging a captured region is harder (prints/syncs change behavior). So: graphs for the steady-state inner loop, eager mode for everything else.

Q6 (leadership): What runtime-layer guarantees would you write into your platform's internal API contract? A: (1) No driver alloc/free in steady state — all hot-path memory from pools; CI fails on cudaMalloc in request paths. (2) Stream-ordered lifetime — the only free is an event-ordered free; raw frees don't exist in the API. (3) Pool segregation + metrics — weights/activations/KV in separate pools, each exporting reserved/allocated/largest-block (the fragmentation dashboard). (4) No device-wide syncs in engine code — event edges only; a lint catches deviceSynchronize. (5) Graph-compatible inner loops — steady-state paths must be capturable (no syncs, no allocs), which doubles as a cleanliness forcing-function. (6) Determinism escape hatch — a debug mode with serialized streams and poisoned frees for reproducing the "garbage tokens once" class. Each guarantee exists because its absence is a specific incident I can name — that's how runtime contracts should be justified.

References

• PyTorch CUDA Caching Allocator notes (the design doc) — https://pytorch.org/docs/stable/notes/cuda.html#cuda-memory-management
• PyTorch CUDACachingAllocator.cpp (read the real one after the lab) — https://github.com/pytorch/pytorch/blob/main/c10/cuda/CUDACachingAllocator.cpp
• NVIDIA, "Using the NVIDIA CUDA Stream-Ordered Memory Allocator" (cudaMallocAsync) — https://developer.nvidia.com/blog/using-cuda-stream-ordered-memory-allocator-part-1/
• NVIDIA, CUDA Graphs section of the Programming Guide + "Getting Started with CUDA Graphs" — https://developer.nvidia.com/blog/cuda-graphs/
• NVIDIA VMM API ("Introducing Low-Level GPU Virtual Memory Management") — https://developer.nvidia.com/blog/introducing-low-level-gpu-virtual-memory-management/
• Wilson et al., "Dynamic Storage Allocation: A Survey and Critical Review" (1995) — the allocator-design bible
• Bonwick, "The Slab Allocator" (USENIX 1994)
• vLLM source: cuda_graph usage and block manager — https://github.com/vllm-project/vllm
• Cross-track: Phase 07 WARMUP (PagedAttention as allocator design), LLM Inference track Phase 09

🛸 Hitchhiker's Guide — Phase 05: Runtime Systems

Read this if: you call torch.cuda.empty_cache() superstitiously, you've never asked why nvidia-smi and your framework disagree, or "stream-ordered deallocation" sounds like jargon. Field notes; derivations in the WARMUP.

0. The 30-second mental model

The runtime layer = pools + DAG. Pools because the driver's alloc/free are slow, synchronizing, and fragmentation is permanent (no paging, no compaction in HBM). DAG because the GPU is async: streams are FIFOs, events are edges, and everything — including freeing memory — must be positioned in that DAG, not in host wall-clock time. Every framework (PyTorch, vLLM, TensorRT) is a specific set of answers to those two problems.

1. The allocator card

• Design: size-binned free lists; round sizes (512B+); split big cached blocks; never return to driver in steady state; OOM → flush cache → retry.
• reserved (nvidia-smi, driver's view) vs allocated (framework's view): the gap is cache, not a leak. Watch largest-free-block for fragmentation.
• Fragmentation defenses (no compaction possible): size bucketing, pool segregation (weights/activations/KV), fixed-block paging for the dominant consumer (= PagedAttention), and — the confession — scheduled restarts.
• cudaMallocAsync = this design, inside the driver, stream-ordered. expandable_segments = VMM tricks to grow reservations.

2. The stream card

• Same stream = ordered. Different streams = concurrent until an event edge says otherwise. Legacy default stream = a global serializer; ban it.
• Review method: draw the intended DAG; extra edges = lost overlap, missing edges = race. Most stream bugs are visible in 5 minutes this way.
• Free is an event: host-time free + async GPU = use-after-free that reproduces once a week. Stream-ordered free (or stream-tagged pools) makes the class impossible, not unlikely.
• CUDA Graphs: capture the decode step, replay in one call; 10–30% at small batch; frozen shapes/pointers; no allocs/syncs inside capture.

3. Numbers

4. Incident patterns

• Day-3 OOM at 60%: checkerboard fragmentation. Look at largest-free-block trend; bucket sizes; segregate pools. Restarts are the tourniquet, not the fix.
• Throughput sag after a "harmless" logging change: someone synced (D2H of a metric) inside the decode loop → killed overlap. nsys timeline shows the gap.
• Garbage tokens, unreproducible: cross-stream use-after-free. Audit every free for stream order; poison freed blocks in a debug build.
• "empty_cache() fixed it": it didn't — it papered over fragmentation by returning blocks to the driver (with syncs). Treat as a smell, not a tool.
• Graph capture failing after a refactor: someone added an alloc/sync in the inner loop. The capture failure is useful — it's a cleanliness linter.

5. What changes at head-of-engineering level

You stop reviewing allocator PRs and start writing the runtime contract: no driver alloc/free in request paths (CI-enforced), stream-ordered lifetime only, segregated pools with reserved/allocated/largest-block metrics exported, no device-wide syncs in engine code, inner loops capture-clean. Each rule maps to an incident class; that mapping is how you defend the contract to a team that finds it pedantic — until day 3.

Next: Phase 06 — with memory and execution under control inside one process, who gets the GPU between processes, pods, and tenants?

Lab 01 — Caching GPU Allocator (Rust)

Phase: 05 — Runtime Systems | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: Rust (safe, stdlib only) | Hardware: none — a CPU-side model of the GPU original

Concept primer: ../WARMUP.md Ch. 2–5.

Run

cargo test                              # 12 correctness tests
cargo run --release --bin fragdemo      # the two teaching demos

0. The mission

Model the two layers every framework stacks on the driver:

• Device = the "driver": a fixed address space, first-fit allocation, holes coalesce on free, and it counts free calls (each is a device sync in real CUDA). This is cudaMalloc/cudaFree.
• CachingAllocator = the framework layer (the PyTorch design): size rounding into classes, binned free lists, block splitting, stream tagging, and flush-and-retry on OOM. This is what actually serves your tensors.

The tests prove the semantics; fragdemo proves why the design exists.

1. What fragdemo shows

== Demo 1: external fragmentation is a shape problem ==
  interleaved (mixed)    free=128 MB  largest hole=  8 MB  16MB alloc: FAILS (fragmented!)
  grouped (disciplined)  free=128 MB  largest hole=128 MB  16MB alloc: OK

== Demo 2: the caching allocator's reason to exist ==
  raw driver :  5000 mallocs,  5000 frees (each free = device sync)
  caching    :     1 device mallocs,  4999 bin hits, 0 device frees

Demo 1 is the day-3 OOM in 30 deterministic lines: the same 128 MB is free in both rows; only the layout differs. Mixed shapes strand free memory in 8 MB holes between permanent blocks, so a 16 MB request fails at 50% utilization — and since HBM can't be compacted (live pointers), the only fix in production is a restart. Shape discipline (grouping lifetimes) keeps one big contiguous hole. This is exactly why Phase 07 pages the KV cache.

Demo 2 is why the caching layer exists at all: 5000 alloc/free cycles become 5000 synchronizing driver frees raw, versus 1 device malloc + 4999 ~100 ns bin hits cached. The driver never sees a steady-state free.

2. Reading order (src/lib.rs)

1. round_size + the Device model — the constraints (first-fit, coalescing, counted frees).
2. CachingAllocator::alloc — the four-step path: bin hit → split larger → driver malloc → flush-and-retry. Each step maps to a WARMUP Ch. 4 bullet.
3. free (goes to a bin, never the driver) and flush_cache (the empty_cache() semantics, including why a split block isn't returnable).
4. The stream-tag logic in alloc — same-stream reuse is free; cross-stream reuse counts an event edge (cross_stream_syncs), the Ch. 7 safety rule.

Then trace a 1.2 MB allocation by hand and check yourself against test_split_accounting.

3. The tests, and what each proves

4. Extensions

1. Coalescing (hard): when two adjacent free blocks exist, merge them back into one larger class — true defragmentation within the cache. Add a test that an alloc/free/alloc pattern which currently fails now succeeds.
2. empty_cache policy: return only segments idle for N cycles; measure the reserved-memory sawtooth.
3. Trace replayer: parse a CSV of (op, size, stream) lines and replay it, emitting reserved/allocated/largest-block per step — turn the allocator into a measurable system you can feed real PyTorch traces.
4. Best-fit vs first-fit in the bins: measure fragmentation on the demo-1 workload; quantify the classic tradeoff.
5. Prometheus metrics: export reserved/allocated/cached/largest-block — the Phase 10 fragmentation dashboard, fed from here.

5. Common pitfalls

1. Returning a split remainder to the driver — you can't free half a segment (test_split_block_not_returnable enforces it; PyTorch has the same limit).
2. Forgetting the stream tag on reuse — silent cross-stream UAF in the real thing; here it's a counted edge so you can see it.
3. Treating reserved as a leak — it's cache; the leak signal is allocated trending up, not reserved.

6. What this lab proves about you

You can implement, test, and reason about the allocator at the heart of every ML framework — and translate its behavior into operational language (reserved vs allocated, fragmentation early-warning, the restart confession). Reading PyTorch's real CUDACachingAllocator.cpp is now a code-review, not an expedition.

Lab 02 — Stream & Event Scheduler (C++)

Phase: 05 — Runtime Systems | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: C++17 (std::thread, no deps) | Hardware: none — a CPU model of CUDA stream semantics

Concept primer: ../WARMUP.md Ch. 6–8.

Run

make && ./scheduler_test

0. The mission

Build the CUDA execution model in ~200 lines and use it to make three production failure modes tangible:

• Stream = a worker thread + FIFO queue. Ops on one stream run in order.
• Event = a counter + condition variable. record(stream) enqueues a bump; stream_wait(s, e) makes stream s block until the event fires. Events are the edges of the dependency DAG (WARMUP Ch. 6).
• Scheduler owns streams and lets the host join the DAG.

Each "op" is just a std::function with a simulated duration — enough to observe overlap, serialization, and lifetime bugs without a GPU.

1. The three scenarios (the heart of the lab)

Expected output:

[scenario 1] serial=282ms  pipelined=140ms  speedup=2.01x   PASS (>=1.7x)
[scenario 2] one-stream=83ms  two-stream=45ms  overlap=1.84x  PASS
             (legacy default stream forces the one-stream case — that's the 2x you lose)
[scenario 3] host-ordered free: CORRUPTION detected (as expected)
             stream-ordered free: no corruption across 200 trials  PASS
[scenario 4] acyclic accepted / circular wait detected   PASS

2. Why each scenario matters

• Scenario 1 is the entire reason streams exist: separate copy and compute engines, wired with event edges, double-buffer so chunk N+1's copy overlaps chunk N's compute (WARMUP Ch. 9). The speedup is the Phase 02 Ch. 9 pipeline, measured.
• Scenario 2 is the bug hiding in every naive codebase: issue to the default stream and your carefully-separated work silently serializes. The fix (per-thread default stream / explicit streams) is a one-line policy with a 2× consequence.
• Scenario 3 is the unreproducible-corruption ticket: the host runs ahead of the device, frees a buffer still in use, the allocator hands it out, and a scheduling coincidence corrupts data once a week. The lab makes it reproducible — then shows stream-ordered free making the class impossible (WARMUP Ch. 7). This connects directly to Lab 01's cross_stream_syncs.

3. Reading order (scheduler.hpp)

1. Stream — worker loop pulling from a std::queue under a mutex/condvar. In-order execution falls out of the FIFO for free.
2. Event — record enqueues a counter bump onto a stream; stream_wait enqueues a blocking wait onto another stream. This is the whole ordering algebra.
3. Scheduler::synchronize — the host joining the DAG (the blunt cudaDeviceSynchronize).
4. The scenarios in scheduler_test.cpp — predict each result from the DAG before reading the assert.

4. Extensions

1. Capture/replay ("mini CUDA Graphs"): record a scenario-1 sequence into a Graph object, then replay it with one call; measure per-op issue cost before/after (WARMUP Ch. 8). Note what you can't capture (data-dependent branches) — the real restriction.
2. Priority streams: a high-priority stream's ops jump the worker queue; show a latency-sensitive op finishing ahead of bulk work.
3. Deadlock detector: build the wait-for graph from event edges and reject cycles at issue time (scenario 4 stubs this).
4. Integrate Lab 01: make the allocator's free stream-ordered using this scheduler's events — the two halves of the runtime layer, joined.

5. Common pitfalls

1. Data race in the test harness itself — the "corruption" in scenario 3 must come from the modeled lifetime bug, not from unsynchronized std::thread access in your scaffolding. The provided harness uses a deterministic interleaving hook; keep it.
2. Busy-waiting instead of condvars — works but pegs a core and hides ordering bugs; use the condition variables as written.
3. Joining workers in the wrong order at shutdown — a classic; the destructor drains queues then joins.
4. Concluding overlap from wall-clock without warmup — same discipline as Phase 02 Ch. 9; the harness discards the first run.

6. What this lab proves about you

You can implement and reason about the execution model that every async GPU program obeys — and you can name, reproduce, and prevent the three failure modes (lost overlap, accidental serialization, cross-stream UAF) that cost real serving fleets latency and correctness. Combined with Lab 01, you've built the runtime layer this JD calls "runtime systems programming."

Phase 06 — GPU Scheduling, Sharing & Virtualization

Difficulty: ⭐⭐⭐⭐☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], Platform/Infra Engineer, Cluster Scheduler Engineer Hardware needed: none — both labs are simulators/models (Python + Go)

Why This Phase Exists

The JD lists "GPU scheduling" as a core skill and the whole role is about building an "orchestration platform." Phase 04 ended on the load-bearing fact: there is no GPU cgroup controller — the kernel cannot enforce a fraction of a GPU. That single gap creates an entire industry: MPS, MIG, time-slicing, fractional-GPU vendors, Kubernetes device plugins, and cluster schedulers that bin-pack the most expensive compute in the datacenter.

This phase is where you learn to allocate GPUs at two altitudes: within a node (how do two tenants share one H100 safely and fairly?) and across a cluster (which of 1,000 GPUs gets this job, and how do you avoid stranding capacity?). You'll build a cluster scheduler simulator that exposes bin-packing, gang scheduling, fragmentation, and fairness, and a Kubernetes device-plugin model that shows exactly how a GPU becomes an allocatable resource — including a fractional-GPU scheme.

Concepts

• The sharing problem restated: no GPU cgroup → sharing needs GPU-side mechanisms
• MPS (Multi-Process Service): one context, multiple processes co-resident; cooperative, weak isolation, no hard memory limits (until recent partitioning)
• MIG (Multi-Instance GPU): hardware partitioning into isolated instances (separate SMs, L2 slices, memory); strong isolation, fixed profiles (1g.10gb …), the only real "fractional GPU"
• Time-slicing: the scheduler context-switches whole-GPU among clients; no isolation, oversubscription convenience
• Choosing among them: latency-sensitive vs batch, trusted vs multi-tenant, the isolation/utilization tradeoff
• Cluster scheduling: bin-packing (first-fit/best-fit/dominant-resource), gang scheduling (all-or-nothing for multi-GPU jobs — the deadlock it prevents), topology-aware placement (NVLink islands, Phase 08)
• Fragmentation at cluster scale: stranded GPUs (a node with 1 free GPU can't take an 8-GPU job); bin-packing policy as the lever
• Fairness: FIFO, priority, preemption, fair-share/DRF; the latency-vs-utilization-vs-fairness trilemma
• Kubernetes integration: the device plugin API, extended resources (nvidia.com/gpu), nvidia-device-plugin, MIG/time-slicing exposure, the scheduler's view
• Inference-specific scheduling: how serving (Phase 07) admission control differs from training job scheduling (Phase 08)

Labs

Lab 01 — GPU Cluster Scheduler Simulator (Python)

Lab 02 — Kubernetes Device Plugin + Fractional GPU (Go)

Deliverables Checklist

• Scheduler sim: policy comparison table reproduced; stranded-GPU and gang results explained
• You can state when to choose MPS vs MIG vs time-slicing from requirements
• Device plugin: full registration→allocate flow passing; fractional advertisement working
• One page: "our platform's GPU-sharing strategy" — which mechanism for which tenant tier (feeds Phase 12)
• You can draw the k8s GPU allocation path (plugin → kubelet → scheduler → pod) from memory

Interview Relevance

• "Compare MPS, MIG, and time-slicing. When each?"
• "An 8-GPU training job won't schedule though 12 GPUs are free. Why, and fix it."
• "What is gang scheduling and what failure does it prevent?"
• "How does Kubernetes know a node has GPUs? Walk me through the device plugin."
• "Your customer wants '0.5 GPU' tiers. How do you actually deliver that, and what do you tell them about isolation?"

Warmup Guide — GPU Scheduling, Sharing & Virtualization

Zero-to-expert primer for Phase 06. Builds on Phase 04 (no GPU cgroup controller) and Phase 01 (SMs, MIG-able hardware). By the end you can choose sharing mechanisms from requirements and design a cluster scheduler.

Table of Contents

• Chapter 1: The Sharing Problem, Restated
• Chapter 2: MPS — Cooperative Co-Residency
• Chapter 3: MIG — Hardware Partitioning
• Chapter 4: Time-Slicing — Oversubscription Without Isolation
• Chapter 5: Choosing a Sharing Mechanism
• Chapter 6: Cluster Scheduling — Bin-Packing and Fragmentation
• Chapter 7: Gang Scheduling and Topology Awareness
• Chapter 8: Kubernetes Device Plugins — How a GPU Becomes Schedulable
• Chapter 9: Fairness, Priority, and Preemption
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: The Sharing Problem, Restated

Phase 04 Ch. 8 established the fact that organizes this entire phase: the Linux kernel has no GPU resource controller. You can cap a process's CPU and memory with cgroups; you cannot cap its share of a GPU's SMs or HBM through any kernel mechanism. The GPU is opaque to the OS scheduler — the driver (and the GPU's own front-end) owns GPU time, and they don't expose a fractional knob to the kernel.

So why share at all? Economics. An H100 costs more than the engineer using it, and a single inference model or notebook frequently uses 10–30% of one. Leaving 70% idle across a fleet is the largest controllable line item in a GPU cloud's P&L. Every mechanism in this phase exists to reclaim that idle capacity without letting tenants corrupt or starve each other.

The three answers, previewed: MPS (let processes co-reside in one context), MIG (cut the silicon into isolated pieces), time-slicing (rotate the whole GPU among clients). They sit at different points on the isolation-vs-utilization curve, and choosing among them per workload is a core platform-leadership skill (Ch. 5).

Chapter 2: MPS — Cooperative Co-Residency

The mechanism: normally, processes sharing a GPU time-slice at context granularity (the GPU switches between their contexts, serializing). The Multi-Process Service runs a daemon that funnels multiple processes' kernels into one shared context, so their kernels execute concurrently on the SMs — filling the gaps a single under-utilizing process leaves.

• Win: higher utilization for many small co-operating jobs (classic case: MPI ranks, or several small inference models). No context-switch overhead between them.
• Isolation: weak. Originally no memory limit (one process could OOM the others' allocations) and no real fault isolation — a fault in one client can poison the shared context (Phase 02 Ch. 4), taking down co-tenants. Volta+ MPS added per-client memory limits and some execution-resource provisioning (CUDA_MPS_ACTIVE_THREAD_PERCENTAGE), narrowing but not closing the gap.
• Verdict: great for trusted, cooperating workloads (your own microservices); not a multi-tenant security boundary.

Chapter 3: MIG — Hardware Partitioning

The mechanism (Ampere+, A100/H100): physically partition one GPU into up to 7 instances, each with dedicated SMs, dedicated L2 cache slices, dedicated memory controllers and HBM, and a separate fault domain. A MIG instance presents to software as its own little GPU.

• Profiles: fixed geometries named <compute>g.<memory>gb — e.g., 1g.10gb (1/7 compute, 10 GB), 2g.20gb, 3g.40gb, 7g.80gb (whole). Memory and compute are partitioned together; you choose from the menu, you don't get arbitrary fractions.
• Isolation: strong — hardware-enforced. A tenant in 1g.10gb cannot touch another instance's memory, cannot steal its SMs, and a fault is contained. This is the only "fractional GPU" with real multi-tenant isolation.
• Cost: rigidity (fixed profiles, reconfiguration usually needs the GPU idle/reset), and quantization waste (a job needing "1.5 instances" rounds up).
• Verdict: the default for multi-tenant fractional serving where isolation is a contractual requirement (sovereign/regulated customers — this JD's nice-to-haves).

Chapter 4: Time-Slicing — Oversubscription Without Isolation

The mechanism: the simplest scheme — the GPU scheduler context-switches the whole GPU among clients in turn (the default behavior when multiple processes use one GPU without MPS/MIG; in k8s, exposed by advertising N "replicas" of one physical GPU).

• Win: trivial oversubscription — pack 4 bursty notebooks onto 1 GPU; when one is idle, others get full GPU during their slice. Zero hardware requirement.
• Isolation: none — clients share all memory (sum must fit!), no performance isolation (a heavy client starves others during its slices), full fault sharing.
• Verdict: dev/test, bursty interactive workloads, trusted oversubscription where occasional interference is acceptable; never for latency-SLO or multi-tenant security.

Chapter 5: Choosing a Sharing Mechanism

The decision table to internalize (and the heart of the "0.5 GPU tier" interview question):

The underlying axis is isolation vs utilization: MIG buys isolation by giving up the flexibility to overcommit; time-slicing buys utilization by giving up isolation; MPS sits between. A platform usually offers tiers mapping to these — and the honest answer to a customer asking for "half a GPU" is "which do you need: a guaranteed isolated half (MIG 3g.40gb), or a cheaper shared half that may be slow when neighbors are busy (time-slicing)?" Saying that sentence is what separates a platform leader from a salesperson.

Chapter 6: Cluster Scheduling — Bin-Packing and Fragmentation

Zoom out from one node to thousands. A scheduler maps jobs (each wanting k GPUs, maybe with constraints) onto nodes (each with some free GPUs). This is bin-packing, NP-hard, solved with heuristics:

• First-fit: place a job on the first node that fits. Fast, simple, packs loosely → strands capacity.
• Best-fit: place on the node that leaves the least leftover. Packs tighter, reduces fragmentation, slightly more expensive.
• Worst-fit / spread: place on the emptiest node. Good for thermal/failure spread and leaving room for big jobs, worse for packing.
• Dominant Resource Fairness (DRF) when jobs want heterogeneous resources (GPUs + CPU + RAM): fair-share along each job's dominant resource.

Cluster fragmentation — the stranded-GPU problem (Lab 01 reproduces it): free GPUs scattered one-per-node can't satisfy an 8-GPU gang job, even though total free ≥ 8. This is Phase 05's external fragmentation at datacenter scale — and the fix is the same family: placement policy (best-fit consolidates; defragmentation by migration where checkpointable), and shape-aware admission (reserve whole nodes for big jobs). The metric to watch: stranded GPUs = free GPUs that no pending job can use given placement constraints.

Chapter 7: Gang Scheduling and Topology Awareness

Gang scheduling: a multi-GPU job needs all its GPUs simultaneously (a distributed training step can't start with 6 of 8 ranks — the all-reduce would hang, Phase 08). So the scheduler must allocate all-or-nothing: either place the whole gang or none of it. Why it matters: without it, two 8-GPU jobs on a 16-GPU cluster can each grab 4–6 GPUs, neither can proceed, neither will release — deadlock. Gang scheduling (and its cousin, reservation) prevents the partial-allocation deadlock. Lab 01 demonstrates both the deadlock and its prevention.

Topology awareness: not all GPU slots are equal. 8 GPUs on one node behind NVLink (900 GB/s, Phase 01 Ch. 9) crush 8 GPUs spread across nodes on InfiniBand (~50 GB/s) for a tensor-parallel job. A topology-aware scheduler prefers NVLink islands for tightly-coupled jobs and is willing to wait or consolidate to get them. This is where Phase 06 (scheduling) and Phase 08 (collectives) meet: the scheduler's placement decides the collective's bandwidth, which decides the job's throughput. Lab 01's extension models NVLink islands with an inter-island penalty.

Chapter 8: Kubernetes Device Plugins — How a GPU Becomes Schedulable

Kubernetes knows nothing about GPUs natively. The device plugin API (a gRPC contract) teaches it. The flow (Lab 02 implements it end to end):

1. The plugin (e.g., nvidia-device-plugin, a DaemonSet pod) registers with the kubelet over a Unix socket, claiming a resource name: nvidia.com/gpu.
2. ListAndWatch: the plugin streams the list of healthy device IDs to the kubelet; the node now advertises nvidia.com/gpu: 8 as an extended resource.
3. The scheduler treats nvidia.com/gpu like CPU/memory for placement (pods request it in resources.limits).
4. When a pod lands, the kubelet calls the plugin's Allocate with the chosen device IDs; the plugin returns the device nodes, mounts, and environment to inject (NVIDIA_VISIBLE_DEVICES=...) — which is exactly the container-injection mechanism from Phase 04 Ch. 9.

The crucial boundary (say it in interviews): the device plugin does allocation, not isolation. It advertises and assigns; the actual isolation is whatever the underlying mechanism provides (whole device, MIG instance, or — for time-slicing — nothing, you just advertised replicas: 4 of one physical GPU). Fractional GPUs in k8s are this: the plugin lists one physical GPU as N allocatable units. MIG appears as distinct resource names (nvidia.com/mig-1g.10gb). Lab 02 builds the plugin, then the fractional advertisement, so the allocation/isolation line is concrete, not slogan.

Chapter 9: Fairness, Priority, and Preemption

The trilemma every scheduler navigates: latency vs utilization vs fairness — you can optimize any two.

• FIFO: simple, starves big/low-priority jobs behind a long queue.
• Priority + preemption: high-priority jobs evict lower ones — but preempting a GPU job means checkpointing (expensive) or killing (wasteful); the cost of preemption is much higher than for CPU jobs, so policies lean toward priority admission over mid-flight eviction.
• Fair-share / DRF: each tenant gets a fair slice over time; prevents one team monopolizing the cluster.
• Backfill: let small jobs run in the gaps while a big gang job waits for its reservation — utilization without starving the big job.

For inference serving specifically (Phase 07), "scheduling" shifts meaning: it's admission control and continuous batching at millisecond granularity on an already-placed model, not job placement at minute granularity. A platform leader keeps these two scheduling worlds — training-job placement (this phase) and request admission (Phase 07) — distinct in design and in conversation.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02 (cluster view, then the node-level k8s mechanism).

Lab 01 (cluster sim, Python):

1. python solution.py; read the policy comparison table, then the scheduler core against Ch. 6–7.
2. Find the stranded-GPU result in first-fit vs best-fit — explain it with the Ch. 6 fragmentation story.
3. Watch the gang test: with naive allocation two big jobs deadlock; with gang scheduling neither does. Connect to Ch. 7.
4. Extensions: preemption (with checkpoint cost), DRF, topology islands.

Lab 02 (device plugin, Go):

1. go test ./... — the mock kubelet drives registration→list→allocate.
2. Read plugin.go against Ch. 8; trace one pod's allocation end to end.
3. Set replicas: 4; re-run; one physical GPU now advertises 4 units. Read the isolation caveat in the test output — that's the allocation≠isolation lesson.
4. Extensions: MIG profiles as resource names, health transitions.

Success Criteria

• You can explain why no GPU cgroup controller forces GPU-side sharing (Ch. 1)
• You can contrast MPS/MIG/time-slicing on isolation, utilization, and use case, and pick from requirements (Ch. 2–5)
• You can explain cluster fragmentation/stranded GPUs and how best-fit helps (Ch. 6)
• You can define gang scheduling and the deadlock it prevents, and the NVLink-island argument (Ch. 7)
• You can draw the k8s device-plugin allocation path and state allocation≠isolation (Ch. 8)
• Both labs pass; you can map sharing mechanisms to platform tiers

Interview Q&A

Q1: Compare MPS, MIG, and time-slicing. A: All three share one physical GPU because the kernel can't. MPS merges multiple processes into one CUDA context so their kernels run concurrently — best utilization for trusted cooperating jobs, but weak isolation (shared fault domain; memory limits only added in Volta+). MIG physically partitions the GPU into instances with dedicated SMs, L2, and HBM and separate fault domains — the only option with hardware-enforced multi-tenant isolation, at the cost of fixed profiles and reconfiguration rigidity. Time-slicing rotates the whole GPU among clients — trivial oversubscription, zero isolation (shared memory, no QoS, shared faults). Decision: trusted+utilization → MPS; multi-tenant+isolation → MIG; bursty dev/cost → time-slicing; latency-SLO → MIG or dedicated.

Q2: An 8-GPU job won't schedule though 12 GPUs are free. Why? A: Cluster fragmentation — the 12 free GPUs are scattered (e.g., 1–2 per node) and the job needs 8 co-located (gang + likely topology constraint for the collective). No single node/island offers 8, so it can't place despite sufficient total capacity — the stranded-GPU problem, external fragmentation at cluster scale. Fixes: best-fit/consolidating placement to keep whole nodes free, reserving capacity for large jobs, draining/migrating checkpointable small jobs to defragment, and topology-aware admission that knows an 8-GPU TP job wants one NVLink island. The metric to alert on is stranded-GPU count, not raw free-GPU count.

Q3: What is gang scheduling and what does it prevent? A: All-or-nothing allocation for a multi-GPU job: the scheduler places every GPU the job needs simultaneously or none. It prevents the partial-allocation deadlock — two 8-GPU jobs on 16 GPUs each grabbing 6, neither able to start its synchronous collective (Phase 08), neither releasing. Gang scheduling (with reservation/backfill so the cluster isn't idled while assembling a gang) is mandatory for distributed training; serving (single-GPU or already-placed replicas) usually doesn't need it.

Q4: How does Kubernetes schedule onto GPUs? A: Via the device-plugin API. A per-node plugin registers a resource name (nvidia.com/gpu) with the kubelet, streams the healthy device list via ListAndWatch (the node now advertises that many extended-resource units), and the default scheduler treats it like CPU/memory for placement. On binding, the kubelet calls the plugin's Allocate with chosen device IDs; the plugin returns the device nodes, mounts, and env (NVIDIA_VISIBLE_DEVICES) to inject — the Phase 04 container-injection mechanism. Critically the plugin does allocation, not isolation: fractional GPUs are just advertising one device as N units (time-slicing, no isolation) or exposing MIG instances as separate resource names (real isolation).

Q5 (leadership): Design the GPU-sharing strategy for a multi-tenant inference platform with both enterprise and free-tier customers. A: Tier by isolation need. Enterprise/regulated: dedicated GPUs or MIG instances (3g.40gb-class) — hardware isolation is contractual; advertise as distinct k8s resources; predictable QoS. Standard paid: MIG smaller profiles or MPS within a trusted pool with per-client memory limits and SM provisioning, SLO-backed. Free/dev tier: time-sliced oversubscription — cheap, explicitly best-effort, isolated only at the container level (and never co-scheduling free tenants with paid on the same physical GPU, to bound interference and the shared-fault blast radius). Across the cluster: best-fit placement to limit stranding, gang+topology scheduling for any training jobs, and a fragmentation dashboard (stranded GPUs, MIG-profile mix utilization). The platform's job is to expose these as a small menu of honest tiers — each customer picks their point on the isolation/cost curve, and we never silently put them somewhere weaker than they paid for. (This is also where Phase 11's tenant-isolation security story attaches.)

References

• NVIDIA Multi-Process Service docs — https://docs.nvidia.com/deploy/mps/
• NVIDIA MIG User Guide — https://docs.nvidia.com/datacenter/tesla/mig-user-guide/
• NVIDIA GPU time-slicing in Kubernetes — https://docs.nvidia.com/datacenter/cloud-native/gpu-operator/latest/gpu-sharing.html
• Kubernetes Device Plugin API — https://kubernetes.io/docs/concepts/extend-kubernetes/compute-storage-net/device-plugins/
• NVIDIA k8s device plugin (read the real one) — https://github.com/NVIDIA/k8s-device-plugin
• Ghodsi et al., "Dominant Resource Fairness" (NSDI 2011)
• "Gang Scheduling" — Feitelson & Rudolph; and the Kubernetes "coscheduling" plugin / Volcano — https://volcano.sh
• Google Borg paper (Verma et al., EuroSys 2015) — the cluster-scheduling bible
• Cross-track: Phase 04 WARMUP Ch. 8–9, Phase 08 WARMUP (why topology matters)

🛸 Hitchhiker's Guide — Phase 06: GPU Scheduling & Virtualization

Read this if: you can run a CUDA job but "how do two tenants share an H100" and "why won't my 8-GPU job schedule" are mysteries. Field notes; derivations in the WARMUP.

0. The 30-second mental model

No GPU cgroup controller ⇒ the OS can't enforce GPU fractions ⇒ sharing needs GPU-side mechanisms (MPS/MIG/time-slicing) and cluster-side scheduling (bin-packing + gang + topology). Every choice is a point on isolation ⇄ utilization. Kubernetes learns about GPUs only through a device plugin, which does allocation, not isolation.

1. Node-level sharing card

"Half a GPU?" → "isolated half (MIG 3g.40gb) or cheap shared half (time-slice that's slow when neighbors are busy)?" Knowing to ask that is the whole skill.

2. Cluster-level card

• Bin-packing: first-fit (fast, strands), best-fit (tight, less stranding), spread (failure/thermal). Stranded GPUs = free GPUs no pending job can use.
• Gang scheduling: all-or-nothing for multi-GPU jobs → prevents the partial-allocation deadlock (two 8-GPU jobs each grabbing 6 on 16 GPUs).
• Topology: 8 GPUs on one NVLink island ≫ 8 spread over IB for TP jobs (Phase 08). Topology-aware placement decides collective bandwidth.
• Backfill + reservation: small jobs fill gaps while a big gang assembles — utilization without starving the big job.

3. The k8s GPU path

device-plugin (DaemonSet) --register--> kubelet --advertises--> nvidia.com/gpu: 8
                          --ListAndWatch (healthy IDs)-->
scheduler places pod requesting nvidia.com/gpu: 1
kubelet --Allocate(ids)--> plugin returns device nodes + mounts + NVIDIA_VISIBLE_DEVICES
                                                       (= Phase 04 injection)

Fractional GPU = plugin advertises one physical device as N units (time-slice, no isolation) OR exposes MIG instances as nvidia.com/mig-1g.10gb (real isolation). Allocation ≠ isolation — tattoo it.

4. Numbers & facts

5. War stories

• "We bought 0.5-GPU tiers and customers complained of jitter" — sold time-slicing as if it were isolated. Free tenants interfered. Fix: MIG for paid isolation tiers; never co-place free + paid on one physical GPU.
• "Cluster at 60% but jobs queueing" — stranded GPUs from first-fit + big-gang jobs. Switched to best-fit + reserved whole-node pool for ≥4-GPU jobs. Queue drained.
• "Training job hangs at start, no error" — partial gang allocation; 6 of 8 ranks up, NCCL init waiting forever. Gang scheduling fixed it.
• "TP throughput half of expected" — scheduler spread 8 GPUs across 2 nodes; per-layer all-reduce hit IB not NVLink. Topology constraint added.
• "MPS client crash took down the node's other models" — shared context, shared fault domain. Moved tenants to MIG.

6. Exit bar

You can pick a sharing mechanism from requirements, explain stranded-GPU fragmentation and gang/topology scheduling, draw the k8s device-plugin path, and state allocation≠isolation with conviction. You've built a cluster scheduler sim and a device plugin. Next: Phase 07 — the most commercially valuable scheduling of all, request scheduling inside an LLM server (KV cache, continuous batching), the JD's named skill.

Lab 01 — GPU Cluster Scheduler Simulator (Python)

Phase: 06 — GPU Scheduling | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: Python (stdlib) | Hardware: none

Concept primer: ../WARMUP.md Ch. 6–7, 9.

Run

python solution.py

0. The mission

Simulate a GPU cluster and run a job trace through four placement policies, measuring what platform leaders actually argue about: utilization, queue wait, stranded GPUs, and gang-deadlock avoidance.

The simulator is event-driven: jobs arrive, request k GPUs for a duration, get placed (or queued), run, and release. Policies differ only in where they place.

Expected output:

== Placement-policy comparison (mixed trace + whole-node jobs) ==
policy            util%  p50_wait  p99_wait  stranded  deadlocks
spread             60.9       0.0     228.0        24         0
first_fit          67.0       0.0      36.0          2         0
best_fit           67.0       0.0      36.0          2         0

== Gang vs naive: the partial-allocation deadlock ==
  naive_no_gang : 0/3 jobs completed, 3 deadlocked (held GPUs, never able to finish)
  best_fit_gang : 3/3 jobs completed, 0 deadlocked

1. What each policy proves

• spread vs best_fit: this is the headline. The cluster runs at the same moderate load, but spread scatters 2-GPU jobs across all four nodes, leaving 2 free GPUs per node and no whole node free — so the whole-node NVLink-8 jobs strand (p99 wait jumps to 228, and 24 GPUs sit free while jobs wait). best_fit consolidates the small jobs onto fewer nodes, keeping whole nodes available → the NVLink jobs run almost immediately. Same hardware, same jobs; only the placement shape differs. This is Phase 05's fragmentation at cluster scale, and the reason topology/consolidation-aware placement exists (WARMUP Ch. 6–7).
• The gang-vs-naive deadlock: three 6-GPU jobs on a 12-GPU cluster. Naive round-robin placement gives each job one 4-GPU node; all 12 GPUs are held, every job is stuck at 4/6, none can finish or release — permanent partial-allocation deadlock. Gang scheduling refuses partial commits and runs them one at a time: all three complete (WARMUP Ch. 7).

2. Reading order (solution.py)

1. Cluster / Node — the resource model and free_gpus, stranded_gpus.
2. place_* policy functions — each takes (cluster, job) → placement or None.
3. The gang logic: naive reserves GPUs as it finds them (can deadlock); gang only commits if the whole job fits atomically.
4. The event loop and metrics.

3. Extensions

1. Preemption: high-priority jobs evict low-priority ones; model the checkpoint cost (a fixed penalty added to the preempted job's remaining time). Show priority-inversion avoidance.
2. DRF fairness: jobs request GPUs and CPU; schedule by dominant share; prevent one tenant monopolizing.
3. Topology depth: model 2 nodes × 4-GPU NVLink islands; add an inter-island bandwidth penalty to job duration when a gang spans islands — then watch topology_aware pay off in throughput, not just placement.
4. Backfill: let small jobs run while a big gang waits for its reservation; measure the utilization recovery.

4. Common pitfalls

1. Counting "free GPUs" as "usable GPUs" — the whole stranded-GPU lesson is that they differ once jobs have co-location constraints.
2. Releasing GPUs in the wrong event order (release before the job's end event) — corrupts utilization accounting.
3. Modeling gang scheduling as "retry until it fits" without a reservation — that's just busy-waiting and can still starve big jobs; the lab uses reservation-style atomic commit.

5. What this lab proves about you

You can reason about cluster scheduling with numbers — stranding, gang-deadlock, topology — which is the difference between "we use Kubernetes" and "here's our placement policy, here's the fragmentation it controls, here's why big jobs don't starve." That's the platform-leadership altitude this JD hires for.

Lab 02 — Kubernetes Device Plugin + Fractional GPU (Go)

Phase: 06 — GPU Scheduling | Difficulty: ⭐⭐⭐⭐☆ | Time: 4–6 hours Language: Go (stdlib) | Hardware: none — a model of the device-plugin contract + a mock kubelet

Concept primer: ../WARMUP.md Ch. 8.

Run

go test ./...        # mock kubelet drives registration -> list -> allocate
go run .             # prints a narrated end-to-end allocation walkthrough

0. The mission

Implement the Kubernetes device-plugin contract for a fake GPU — the real mechanism by which a node advertises nvidia.com/gpu: 8 and a pod gets device nodes injected. Then extend it to advertise fractional GPUs (time-sliced replicas) and MIG profiles as distinct resources, so the allocation ≠ isolation boundary (WARMUP Ch. 8) is concrete code, not a slogan.

We model the three plugin operations that matter (the real API is gRPC over a Unix socket; we keep the semantics and the data flow, in plain Go interfaces, so it runs anywhere with no proto toolchain):

1. What the walkthrough shows

mode=whole       resource=nvidia.com/gpu        advertised units=4   (4 physical GPUs)
  allocate [GPU-0 GPU-1] -> env NVIDIA_VISIBLE_DEVICES=GPU-0,GPU-1
                            devices /dev/nvidia0 /dev/nvidia1
  isolation: STRONG (dedicated physical devices)

mode=timeslice   resource=nvidia.com/gpu        advertised units=16  (4 GPUs x 4 replicas)
  allocate [GPU-0::3] -> env NVIDIA_VISIBLE_DEVICES=GPU-0  (replica 3 shares GPU-0!)
  isolation: NONE (replicas share memory + faults — WARMUP Ch. 4)

mode=mig         resource=nvidia.com/mig-1g.10gb advertised units=28  (4 GPUs x 7 instances)
  allocate [MIG-0-2] -> env NVIDIA_VISIBLE_DEVICES=MIG-GPU-0/2/0
  isolation: STRONG (hardware-partitioned instance)

The same plugin API advertises all three — but the isolation each delivers is completely different. Time-slicing's "16 GPUs" are 4 physical devices wearing 4 hats each; the plugin happily allocates replica 3 of GPU-0 to a pod with no memory or fault isolation from replicas 0–2. That's the lesson the lab burns in: the device plugin allocates; it does not isolate.

2. Reading order

1. device.go — the Device, Plugin interface, and the three modes (whole / timeslice / mig) that change only what gets advertised.
2. plugin.go — ListAndWatch (build the advertised list) and Allocate (pick devices, build the injection response with NVIDIA_VISIBLE_DEVICES).
3. kubelet_mock.go — the mock kubelet: registers the plugin, reads the device list, and requests an allocation, exactly as the real kubelet does.
4. plugin_test.go — the contract tests.

3. Extensions

1. MIG geometry: advertise mixed profiles (1g.10gb, 2g.20gb, 3g.40gb) as separate resource names from the same physical GPUs, respecting that a GPU's instances are mutually exclusive (allocating a 3g consumes capacity the 1gs would have used) — the real MIG accounting problem.
2. Health transitions: mark a device Unhealthy (model an Xid error from Phase 04/10); ListAndWatch must re-advertise and the kubelet must stop scheduling onto it.
3. gpu-feature-discovery: emit node labels (nvidia.com/gpu.product=H100, nvidia.com/gpu.memory=80) so the scheduler can do affinity.
4. Preferred allocation: implement GetPreferredAllocation to pick devices that share NVLink (topology — Phase 08) for multi-GPU requests.

4. Common pitfalls

1. Conflating advertised units with isolation — the entire point. A test asserts that two timeslice replicas of the same physical GPU resolve to the same NVIDIA_VISIBLE_DEVICES value (proof they share hardware).
2. Forgetting the env/mount injection — without it the container sees the resource accounted but can't actually reach the GPU (Phase 04 Ch. 9).
3. Letting Allocate hand out more replicas than the device has, or replicas of an unhealthy device.

5. What this lab proves about you

You can implement the exact integration point between Kubernetes and GPUs, and you can articulate — with running code — why "fractional GPU" means very different isolation guarantees depending on the mechanism. When a customer or a PM asks for "GPU sharing in k8s," you can design the honest answer (which resource, which isolation, which tier — Phase 12).

Phase 07 — LLM Serving & KV-Cache Management

Difficulty: ⭐⭐⭐⭐⭐ | Estimated Time: 2.5 weeks Roles supported: Head of Engineering [GPU], LLM Inference Engineer, Serving Platform Lead Hardware needed: none — both labs are CPU models of the serving runtime (Python); optional GPU for the vLLM reading exercise

Why This Phase Exists

The JD names it explicitly: "GPU scheduling, KV-cache management, or LLM serving optimisation" and "deploying and optimising AI infrastructure at production scale, ideally with multiple inference engines." This is the single most commercially valuable phase for the role — a 30% throughput gain across a serving fleet is millions of dollars a year, and the lever is almost always KV-cache management and batching.

This is the consequence phase. Phase 01 proved decode is memory-bandwidth-bound; Phase 05 built allocators and stream scheduling; Phase 06 built request-vs-job scheduling. Here it all converges into the serving runtime: you'll build a paged KV-cache allocator with prefix sharing (PagedAttention as an allocator-design problem — exactly Phase 05 Ch. 5's prediction) and a continuous-batching simulator with TTFT/TPOT/goodput, then connect both to operating vLLM/TensorRT-LLM/SGLang at scale.

This phase deliberately overlaps the LLM Inference Engineer track's Phase 09. That track owns the model/algorithm depth (FlashAttention math, quantization formats, speculative decoding derivations); this phase owns the platform/operations view — the cache as a managed resource, the scheduler as a fleet control loop, multi-engine selection, and the leadership decisions. Read both; they're complementary, and the WARMUPs cross-reference rather than duplicate.

Concepts

• The prefill/decode split and why it dictates everything (recap from Phase 01 Ch. 7, serving-side)
• KV-cache memory math: 2 × n_layers × n_kv_heads × d_head × seq_len × dtype_bytes — and why it, not weights, caps concurrency
• PagedAttention: fixed-size KV blocks + block tables = external fragmentation eliminated for the dominant consumer (Phase 05 Ch. 5, realized)
• Prefix caching / sharing: copy-on-write KV blocks for shared system prompts (chat APIs) — the highest-ROI serving feature for instruction-heavy workloads
• Continuous batching (Orca): iteration-level scheduling, admission/eviction, the latency-throughput dial
• Chunked prefill and prefill/decode disaggregation: keeping long prompts from stalling decodes
• Metrics & SLOs: TTFT, TPOT/ITL, E2E, throughput, and goodput (throughput meeting SLO — the only honest fleet metric)
• Admission control, preemption, and KV-cache swapping/recompute under memory pressure
• GQA, KV quantization (INT8/FP8 KV) — shrinking the cache (cross-ref to llm-inference track)
• The engine landscape as a platform decision: vLLM vs TensorRT-LLM vs SGLang vs TGI — what each optimizes, when to choose which, the multi-engine reality
• Sovereign/air-gapped serving: on-prem deployment, no telemetry phone-home, license enforcement (Phase 11) — the JD's regulated-industry nice-to-haves

Labs

Lab 01 — Paged KV-Cache Allocator + Prefix Sharing (Python)

Lab 02 — Continuous Batching Simulator + Goodput (Python)

Deliverables Checklist

• Paged allocator: tests pass; memory-comparison table reproduced and explained
• You can do the KV-cache memory math cold for any model/context/batch
• Continuous-batching sim: metrics table reproduced; goodput-optimal batch depth identified
• You can explain why goodput > throughput as a fleet metric
• One page: "our serving platform's KV-cache & batching strategy" (feeds Phase 12 capstone)
• vLLM reading: skim vllm/core/block_manager + scheduler; map each to your Lab 01/02 code

Interview Relevance

This phase is the direct portfolio for the JD's serving requirement.

• "Walk me through KV-cache memory and why it caps concurrency."
• "Explain PagedAttention — and why it's really an allocator." (Phase 05 link)
• "Static vs continuous batching — and what's goodput?"
• "How does prefix caching work and when is it worth it?"
• "vLLM vs TensorRT-LLM vs SGLang — when each, and how do you run multiple engines?"
• "Design air-gapped LLM serving for a regulated customer." (→ Phase 11, system-design/)

Warmup Guide — LLM Serving & KV-Cache Management

Zero-to-expert primer for Phase 07, from the platform/operations angle. Builds on Phase 01 (roofline), Phase 05 (allocators), Phase 06 (scheduling). For the model/algorithm depth (FlashAttention math, quantization formats, speculative decoding) see the LLM Inference track Phase 09; this guide cross-references rather than duplicates.

Table of Contents

• Chapter 1: The Serving Problem in One Diagram
• Chapter 2: KV-Cache Memory — the Math That Caps Everything
• Chapter 3: PagedAttention — KV Cache Is an Allocator Problem
• Chapter 4: Prefix Sharing — Copy-on-Write for Prompts
• Chapter 5: Continuous Batching — Iteration-Level Scheduling
• Chapter 6: Metrics and SLOs — Why Goodput Is the Only Honest Number
• Chapter 7: Memory Pressure — Preemption, Swapping, Recompute
• Chapter 8: The Engine Landscape — A Platform Decision
• Chapter 9: Sovereign and Air-Gapped Serving
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: The Serving Problem in One Diagram

One LLM request is two workloads (Phase 01 Ch. 7's roofline, applied):

PREFILL  (process the whole prompt)      DECODE (generate token by token)
  all prompt tokens at once                one token per forward pass
  big matmuls -> COMPUTE-bound             reads all weights+KV -> BANDWIDTH-bound
  cost ~ prompt length                     cost ~ output length
  user feels it as: TTFT                   user feels it as: TPOT (smoothness)

Everything in serving follows from three facts: (1) decode is bandwidth-bound, so batching (reusing each weight read across many requests) is the master throughput lever; (2) the KV cache grows with every token and caps how many requests you can batch; (3) prefill and decode interfere (a long prefill stalls everyone's decode). The platform's job is to maximize goodput — useful throughput meeting latency SLOs — by managing the cache as a resource and scheduling at iteration granularity. Hold that; the rest elaborates.

Chapter 2: KV-Cache Memory — the Math That Caps Everything

The formula to know cold (and derive in interviews):

$$\text{KV bytes} = 2 \times n_{layers} \times n_{kv_heads} \times d_{head} \times \text{seq_len} \times \text{bytes/elem}$$

The 2 is K and V. Worked: Llama-3-8B (32 layers, 8 KV heads — GQA already applied — , d_head 128, BF16): per token = 2·32·8·128·2 = 131 KB/token; at 8K context that's ~1 GB per sequence, on top of the 16 GB of weights. A 40 GB GPU holds ~22 such sequences max — the KV cache, not compute, caps concurrency.

This single calculation explains the whole phase:

• GQA (8 KV heads instead of 32) already cut this 4× — why every modern model uses it.
• KV quantization (INT8/FP8 KV) halves/quarters it again.
• Paging (Ch. 3) stops you from wasting most of it.
• Batching limits (Ch. 5) are set by how many sequences' caches fit.

As a platform leader you carry this formula to capacity planning: "how many concurrent 8K-context users per H100?" is (HBM − weights) / (KV bytes/token × context), and every serving optimization is an attempt to improve one term.

Chapter 3: PagedAttention — KV Cache Is an Allocator Problem

Phase 05 Ch. 5 predicted this exactly: "fixed-size paging for the dominant consumer eliminates external fragmentation by design — that sentence is PagedAttention." Here it is, realized (Lab 01 builds it).

The naive way: allocate each sequence a contiguous KV buffer sized for max possible length. Problems: (1) you don't know the output length, so you reserve for the worst case → most of it is never used (internal fragmentation, vLLM measured 60–80% waste); (2) variable-length sequences leave un-fillable holes (external fragmentation, Phase 05's day-3 OOM).

PagedAttention: chop the KV cache into fixed-size blocks (e.g., 16 tokens). Each sequence gets a block table (logical block → physical block), exactly like OS virtual memory page tables. Allocate blocks on demand as the sequence grows; the attention kernel gathers K/V through the block table (one indirection).

Wins, each a direct allocator consequence:

• No external fragmentation — all blocks identical size, any free block fits any sequence.
• Near-zero internal fragmentation — waste is at most one partial block per sequence (vs reserving max length).
• 2–4× more concurrent sequences in the same HBM → 2–4× throughput, because more sequences = bigger batches = better decode amortization (Ch. 5).
• Sharing becomes trivial — blocks are relocatable units that multiple sequences can reference (Ch. 4).

The one tradeoff: block size. Bigger blocks = smaller block tables, more internal fragmentation; smaller blocks = more table overhead, finer packing. 16 tokens is the common sweet spot. Lab 01 lets you measure the curve.

Chapter 4: Prefix Sharing — Copy-on-Write for Prompts

The highest-ROI feature for real workloads, and a clean payoff of paging:

Chat and agent APIs send the same long system prompt on every request ("You are a helpful assistant... [2000 tokens of instructions]"). Naively, each request recomputes and stores that prompt's KV — wasted compute (prefill) and wasted memory (identical KV copies).

Prefix caching: hash the prompt prefix; if its KV blocks already exist, share them — multiple sequences' block tables point at the same physical blocks. New requests skip the shared prefill entirely (TTFT collapses) and add no memory for the shared part.

The mechanism is copy-on-write (Lab 01 implements it): shared blocks are read-only and ref-counted; when a sequence diverges (its tokens differ from the shared prefix, or it writes into a shared block), that block is copied first. This is fork() semantics for KV cache — and beam search / parallel sampling (N continuations of one prompt) use the same machinery.

ROI: for a chat assistant with a 2K-token system prompt at 1000 req/s, prefix caching can cut prefill compute by >90% and is often the single biggest serving win. As a platform decision: enable it, measure hit rate, and expose prompt structure guidance to customers (stable prefix first) — a product lever, not just a flag.

Chapter 5: Continuous Batching — Iteration-Level Scheduling

Batching is the master throughput lever (decode reads all weights per step regardless of batch size — Phase 01 Ch. 7 — so batch B amortizes that read B ways). But how you batch decides everything.

• Static batching: collect B requests, run them together until all finish. Fails for LLMs: requests finish at wildly different lengths, so the batch runs at the speed of its longest member while finished slots sit idle (head-of-line blocking), and new arrivals wait for the whole batch (batch-formation latency). GPU utilization collapses under length variance.

• Continuous batching (Orca, 2022; every modern engine): schedule at iteration granularity. After each decode step: retire finished sequences (free their KV blocks), admit waiting requests (run their prefill, add to the batch). The batch composition changes every step; utilization stays high under length variance; new requests start almost immediately. Typical win: 3–10× throughput over static. Lab 02 measures it.

The induced problems (which the rest of serving solves):

• Admitting a request commits to its KV growth → memory pressure → preemption/swapping (Ch. 7).
• A long prefill entering the batch stalls every decode that iteration → chunked prefill (slice the prompt across iterations) and prefill/decode disaggregation (separate GPU pools, ship KV between them).
• The latency-throughput dial: deeper batches = more throughput, worse TPOT per user (more work per iteration). You don't maximize batch depth — you maximize goodput (Ch. 6).

Chapter 6: Metrics and SLOs — Why Goodput Is the Only Honest Number

The serving vocabulary, and the one metric that resists gaming:

• TTFT (time to first token) = queue wait + prefill. The "feels responsive" metric; dominated by prompt length and batch interference.
• TPOT / ITL (time per output token) = decode speed. The "streams smoothly" metric; human reading speed ~10–15 tok/s is the UX floor.
• E2E latency = TTFT + TPOT × output_len.
• Throughput = total tokens/sec across the fleet.
• Goodput = throughput that meets SLO. The honest metric, because raw throughput is gameable: you can inflate tokens/sec by running enormous batches that blow every request's TPOT past the SLO — "high throughput," zero happy users. Goodput counts only tokens delivered within their latency target.

Why this is a leadership point: vendors and internal teams report throughput because it's bigger; you require goodput at stated p50/p99 SLOs under a realistic arrival process (Poisson, not back-to-back) with the real prompt/output length distribution. A benchmark missing those caveats is marketing (same discipline as Phase 01 Ch. 10's spec-sheet skepticism). Lab 02 shows goodput peaking at an interior batch depth — proof that "maximize batch" is wrong and "maximize goodput" is right.

Chapter 7: Memory Pressure — Preemption, Swapping, Recompute

Continuous batching admits requests optimistically; sometimes their combined KV growth exceeds HBM. The platform must degrade gracefully:

• Preemption: evict a running sequence to free its blocks. Two recovery options when it's rescheduled:
  • Recompute: throw away its KV, re-run prefill later. Cheap memory, costs compute (re-prefill). Best when prefill is short.
  • Swap: copy its KV blocks to host RAM, copy back later. Costs the 50× PCIe bandwidth cliff (Phase 01 Ch. 9), but preserves the work. Best when the KV is large and prefill was expensive.
• Admission control: the cleaner lever — don't admit a request you can't sustain; queue it. This protects p99 for admitted requests (the goodput argument) at the cost of TTFT for queued ones. SLO-aware admission (shed or queue load to protect tails) is the production posture.

The policy choice (recompute vs swap vs reject) is workload-dependent and is a real config surface you own. vLLM exposes both; knowing why you'd pick each is the Phase-07 depth interviewers probe.

Chapter 8: The Engine Landscape — A Platform Decision

The JD wants "multiple inference engines." The map, as a selection problem:

• vLLM: PagedAttention's home; best-in-class throughput, broad model support, fast-moving OSS, great for heterogeneous/multi-model fleets and rapid adoption of new models. The default starting point.
• TensorRT-LLM: NVIDIA's engine; compiles model-specific optimized kernels (ahead-of-time, per-GPU — Phase 03's AOT story) for peak latency/throughput on NVIDIA hardware — at the cost of build complexity and NVIDIA lock-in. Best when you've fixed your model and hardware and need the last 20%.
• SGLang: strong on structured generation, complex prompting, and RadixAttention (advanced prefix sharing); excellent for agent/tool workloads.
• TGI (HF): solid, integrated with the HF ecosystem; common for simpler deployments.

The platform reality is multi-engine: different models/SLAs/hardware want different engines, and you abstract them behind a common gateway (the capstone). The selection axes: throughput vs latency, model coverage vs peak per-model, build complexity vs portability (and the hardware-lock-in question that Phase 09's HAL exists to manage). "We standardized on one engine" is rarely the right answer at scale; "we have a serving abstraction and route models to the engine that serves them best" is.

Chapter 9: Sovereign and Air-Gapped Serving

The JD's regulated-industry nice-to-haves (defence, finance, healthcare, government) reshape serving:

• No phone-home: no telemetry, no model-download-at-runtime, no license server callout. Everything — weights, engine, license validation (Phase 11) — must work offline. This changes your packaging, your update story, and your observability (local-only, Phase 10).
• On-prem capacity is fixed: no elastic cloud burst, so admission control and goodput discipline (Ch. 6) matter more — you can't scale out of a traffic spike.
• License enforcement without a server: signed, offline-verifiable licenses with hardware binding and grace periods (Phase 11 Lab 02) — because the customer's airtight network can't reach your licensing API.
• Auditability: regulated customers need request/response audit trails, often with data residency guarantees — a serving-layer feature, not an afterthought.

This is where Phases 07, 09 (portability for on-prem hardware diversity), and 11 (licensing/protection) converge into the JD's "sovereign AI" pitch — and it's a genuine market differentiator, because most serving stacks assume the cloud.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02 (the batching sim's admission logic depends on the paged allocator's concepts).

Lab 01 (paged KV allocator):

1. python solution.py; read the allocator against Ch. 3–4 — BlockAllocator, BlockTable, the COW logic.
2. Reproduce the memory comparison: naive contiguous vs paged vs paged+prefix. Tie each number to Ch. 2–3 (internal/external fragmentation eliminated).
3. The prefix-sharing demo: N sequences sharing a system prompt collapse to one physical copy; at divergence, COW kicks in. Trace the ref-counts.
4. Extensions: eviction (recompute vs swap, Ch. 7), then a beam-search fork.

Lab 02 (continuous batching + goodput):

1. python solution.py; read the scheduler against Ch. 5–6.
2. Reproduce static-vs-continuous throughput; then sweep max_batch and find where goodput peaks (interior, not max) — the Ch. 6 thesis, measured.
3. Watch a long prefill stall decodes; reason about chunked prefill (extension).
4. Extension: integrate Lab 01's prefix cache (TTFT collapse on cache hits).

Success Criteria

• You can derive KV-cache bytes and concurrent-sequence capacity for any model/context cold (Ch. 2)
• You can explain PagedAttention as an allocator and the block-size tradeoff (Ch. 3)
• You can explain prefix sharing + COW and when it's the biggest win (Ch. 4)
• You can contrast static vs continuous batching and explain the latency-throughput dial (Ch. 5)
• You can define goodput and argue why it beats throughput as a fleet metric (Ch. 6) — and your Lab 02 result proves it
• You can choose among vLLM/TensorRT-LLM/SGLang from requirements (Ch. 8)
• You can describe what air-gapped serving changes (Ch. 9)

Interview Q&A

Q1: Derive the KV cache size for Llama-3-8B at 8K context and tell me what it implies for serving. A: KV bytes = 2 × layers × kv_heads × d_head × seq × bytes = 2·32·8·128·8192·2 ≈ 1.07 GB per sequence (8 KV heads because of GQA; without it, 32 heads → 4× more). On a 40 GB GPU holding 16 GB of weights, ~22 GB remains → ~20 concurrent 8K sequences max, and that's the hard cap on batch size — not compute. So serving economics are dominated by fitting more sequences: GQA (already in the 8), KV quantization (FP8 → 2×), and paging to stop wasting the budget on max-length reservations. Capacity planning is (HBM − weights) ÷ (KV/token × context).

Q2: Explain PagedAttention, and why you'd call it an allocator rather than an attention algorithm. A: The attention math is unchanged; what changes is KV memory management. Naive serving reserves a contiguous max-length KV buffer per sequence, wasting 60–80% (you don't know the output length) and fragmenting under variable lengths. PagedAttention stores KV in fixed-size blocks (≈16 tokens) with a per-sequence block table mapping logical to physical blocks — OS virtual memory for the KV cache. That eliminates external fragmentation (uniform blocks) and caps internal fragmentation at one partial block per sequence, yielding 2–4× more concurrent sequences and therefore throughput. It's an allocator because the problem it solves is allocation — fixed-size paging for the dominant consumer, exactly the generic fragmentation fix; the attention kernel just gains one indirection through the block table.

Q3: Static vs continuous batching, and what is goodput? A: Static batching forms a fixed batch and runs it to completion — for LLMs that means the batch runs at its longest member's speed while finished sequences waste their slots (head-of-line blocking), and arrivals wait for batch formation; utilization collapses under length variance. Continuous batching schedules at iteration granularity: after every decode step it retires finished sequences and admits new ones, so the batch is always full of useful work — typically 3–10× throughput. Goodput is the throughput that meets the latency SLO; it matters because raw throughput is gameable by oversized batches that blow everyone's TPOT — high tokens/sec, zero satisfied users. You tune the batch depth to maximize goodput (an interior optimum), not throughput.

Q4: A chat product has a 2000-token system prompt on every request. What do you do? A: Prefix caching. Hash the shared prefix and share its KV blocks across all requests via copy-on-write: new requests skip re-prefilling the system prompt (TTFT collapses) and add zero memory for the shared part; blocks are ref-counted and copied only when a sequence diverges. For a 2K-token prompt at high QPS this can cut prefill compute >90% — usually the single biggest win. As a platform move I'd also advise the product to keep the stable instructions first (so the cacheable prefix is maximal) and monitor cache hit rate as a first-class metric. SGLang's RadixAttention generalizes this to tree-structured shared prefixes for agent workloads.

Q5: Under memory pressure mid-generation, what are your options? A: First, admission control — ideally don't admit a request you can't sustain; queue it (protects p99 for admitted requests, the goodput argument). For already- running sequences, preempt and recover one of two ways: recompute (drop the KV, re-prefill on reschedule — cheap memory, costs compute; good when prefill is short) or swap (copy KV to host RAM and back — preserves work but pays the 50× PCIe cliff; good when KV is large and prefill was expensive). The choice is workload-dependent and a config surface I'd own. The anti-pattern is admitting unboundedly and OOM-crashing the engine, which kills every in-flight request.

Q6 (leadership): How do you choose inference engines for a multi-model platform? A: I don't standardize on one. Different models, SLAs, and hardware want different engines: vLLM as the throughput-strong, broad-coverage default and for fast adoption of new models; TensorRT-LLM where a model and GPU are fixed and we need the last 20% of latency/throughput (accepting build complexity and NVIDIA lock-in — the lock-in our Phase 09 HAL exists to bound); SGLang for agent/structured-generation workloads via RadixAttention. The architecture is a common serving gateway abstracting the engines, routing each model to the engine that serves it best, with goodput-at-SLO (not vendor throughput numbers) as the selection evidence. For sovereign customers, add the constraint that the engine must run fully air-gapped with offline licensing — which narrows the field and shapes packaging. The platform value is the abstraction + routing + honest benchmarking, not any single engine.

References

• Kwon et al., "Efficient Memory Management for Large Language Model Serving with PagedAttention" (vLLM, SOSP 2023) — https://arxiv.org/abs/2309.06180
• Yu et al., "Orca: A Distributed Serving System for Transformer-Based Generative Models" (OSDI 2022) — continuous batching
• Zheng et al., "SGLang: Efficient Execution of Structured Language Model Programs" (RadixAttention) — https://arxiv.org/abs/2312.07104
• Agrawal et al., "Sarathi-Serve: Taming Throughput-Latency Tradeoff" (chunked prefill) — https://arxiv.org/abs/2403.02310
• vLLM docs + source (block_manager, scheduler) — https://docs.vllm.ai , https://github.com/vllm-project/vllm
• NVIDIA TensorRT-LLM — https://github.com/NVIDIA/TensorRT-LLM
• DistServe (prefill/decode disaggregation) — https://arxiv.org/abs/2401.09670
• Cross-track: LLM Inference Engineer Phase 09 WARMUP (model/algorithm depth), Phase 05 WARMUP Ch. 5 (paging as allocator design)

🛸 Hitchhiker's Guide — Phase 07: LLM Serving & KV-Cache Management

Read this if: you can run vllm serve but couldn't yet derive its memory limits, explain PagedAttention as an allocator, or defend "goodput, not throughput" to a skeptical exec. Field notes; derivations in the WARMUP and the LLM Inference track Phase 09.

0. The 30-second mental model

Decode is bandwidth-bound ⇒ batching is the throughput lever. The KV cache grows per token and caps batch size ⇒ manage it like memory: page it (PagedAttention), share it (prefix caching), shrink it (GQA, FP8 KV), evict it under pressure. Schedule at iteration granularity (continuous batching). Optimize goodput (throughput meeting SLO), never raw throughput.

1. The one formula

KV bytes = 2 · layers · kv_heads · d_head · seq_len · bytes_per_elem
concurrent_seqs ≈ (HBM − weights) / (KV_bytes_per_token · context)

Llama-3-8B, 8K, BF16 ≈ 1 GB/seq ⇒ ~20 seqs on a 40 GB GPU. The cache, not compute, caps concurrency. Every optimization improves one term.

2. The KV-cache management ladder

PagedAttention = OS virtual memory for the KV cache. Block size ≈16 tokens (table overhead vs internal fragmentation tradeoff).

3. The batching/scheduling ladder

• Static = collect, run to completion. Head-of-line blocking + formation delay. Don't.
• Continuous (Orca) = iteration-level admit/retire. 3–10× throughput. The standard.
• Chunked prefill = slice long prompts so they don't stall decodes.
• Disaggregation = separate prefill and decode GPU pools, ship KV between. Frontier-scale.
• The dial: deeper batch = more throughput, worse per-user TPOT. Tune for goodput, which peaks at an interior batch depth (Lab 02 proves it).

4. Metrics that resist gaming

Require goodput at p50/p99 SLOs, Poisson arrivals, real length distribution. Anything else is marketing (Phase 01 Ch. 10 energy).

5. Engine selection card

Platform answer: a serving gateway abstracting engines, route per model, select on goodput evidence. Not "one engine to rule them all."

6. Air-gapped/sovereign changes

No phone-home (offline weights, offline license — Phase 11), fixed on-prem capacity (admission/goodput matter more, no elastic burst), local-only observability (Phase 10), audit trails + data residency. Most stacks assume cloud; building serving that doesn't is the JD's sovereign-AI differentiator.

7. War stories

• "Throughput great in the benchmark, users furious" — benchmark ran giant batches back-to-back; real Poisson traffic blew p99 TTFT. Switched to goodput-at-SLO reporting; cut batch depth; users happy, "throughput" down 15%.
• "OOM crash takes down 80 in-flight requests" — no admission control, optimistic batching hit the KV wall. Added admission + preempt-with-recompute.
• "Chat TTFT randomly 2s" — no prefix caching; every request re-prefilled a 3K system prompt. Enabled it; TTFT → 80ms on cache hits.
• "New model 40% slower on TRT-LLM than vLLM" — TRT-LLM engine built for the old shape; nobody rebuilt. AOT engines need a rebuild per model/GPU (Phase 03). Routed it to vLLM until rebuilt.

8. Exit bar

You can derive KV memory and capacity cold, explain PagedAttention/prefix caching/continuous batching as the allocator-and-scheduler problems they are, defend goodput over throughput, select engines from requirements, and describe air-gapped serving. Plus you've built the paged allocator and the batching sim. Next: Phase 08 — when one GPU isn't enough, how collectives and topology turn many GPUs into one.

Lab 01 — Paged KV-Cache Allocator + Prefix Sharing (Python)

Phase: 07 — LLM Serving | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: Python (stdlib) | Hardware: none

Concept primer: ../WARMUP.md Ch. 2–4.

Run

python solution.py        # correctness tests + memory comparison + prefix-share demo

0. The mission

Build PagedAttention's memory manager (not the attention math — the allocator, which is the actual innovation, Phase 05 Ch. 5):

• BlockAllocator — a pool of fixed-size KV blocks (16 tokens each).
• BlockTable — a sequence's logical→physical block map (OS page table).
• Copy-on-write prefix sharing — multiple sequences referencing the same physical blocks for a shared system prompt, with ref-counting and copy-on- divergence.

Then measure the three regimes the WARMUP describes.

1. What the output shows

== memory efficiency: sequences that fit in a 4096-block pool ==
naive contiguous (reserve max_len=2048)  :   32 sequences
paged (allocate-on-demand, actual=768)    :   85 sequences   (2.7x)

== prefix sharing: 50 requests, shared 1024-token system prompt ==
without sharing :  3350 blocks used
with sharing    :   214 blocks used   (15.7x less; shared prefix stored once)
copy-on-write   : divergence forked block; refcounts correct

(The prefix-sharing ratio is large here because the shared 1024-token prompt dwarfs each request's 40-token generation — exactly the chat-serving regime where prefix caching pays off most.)

• Naive contiguous reserves max_len per sequence → most blocks unused (internal fragmentation), so few sequences fit. Paged allocates only what each sequence actually uses → 2–4× more concurrent sequences (WARMUP Ch. 3).
• Prefix sharing stores the shared system prompt's blocks once; 50 requests reference them instead of copying — the chat-serving win (WARMUP Ch. 4).
• Copy-on-write forks a shared block only when a sequence diverges, ref-counts staying correct.

2. Reading order (solution.py)

1. BlockAllocator — allocate/free, the free list, ref-counts.
2. BlockTable / Sequence — append tokens, grow blocks on demand.
3. share_prefix + the COW path — ref-count increment on share, copy on write to a shared (refcount>1) block.
4. The three demos in main().

3. Extensions

1. Eviction under pressure: when the pool is full, evict an LRU sequence; implement both recovery policies — recompute (drop KV) and swap (move blocks to a host-side pool) — and compare their cost (WARMUP Ch. 7).
2. Beam search / parallel sampling: fork a sequence into N continuations sharing the prompt's blocks; show COW only on the diverging tokens.
3. Block-size sweep: vary block size (4/8/16/32 tokens); plot internal fragmentation vs block-table size — find the sweet spot (WARMUP Ch. 3).
4. Hash-based prefix cache: auto-detect shared prefixes by hashing block contents (RadixAttention-lite) instead of explicit sharing.

4. Common pitfalls

1. Freeing a shared block while refcount > 1 — corrupts other sequences; free must decrement and only reclaim at zero.
2. Writing into a shared block without COW — silently corrupts every sharer (the test test_cow_on_divergence catches it).
3. Off-by-one at block boundaries — token T lives in block T // block_size at offset T % block_size; the append logic must roll to a new block exactly at the boundary.
4. Counting "blocks used" as allocated-not-shared — the memory win is in physical blocks; shared blocks count once.

5. What this lab proves about you

You've implemented the core of vLLM's memory manager and can explain it as the allocator it is — fixed-size paging + COW sharing for the dominant memory consumer. Reading vllm/core/block_manager_v2.py is now a code review; you can reason about block size, prefix-cache hit rate, and eviction policy as the platform knobs they are.

Lab 02 — Continuous Batching Simulator + Goodput (Python)

Phase: 07 — LLM Serving | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: Python (stdlib) | Hardware: none

Concept primer: ../WARMUP.md Ch. 5–6.

Run

python solution.py

0. The mission

Simulate an LLM server under Poisson request load and prove two things with numbers: continuous batching beats static, and goodput peaks at an interior batch depth (so "maximize batch" is wrong, "maximize goodput" is right — WARMUP Ch. 6).

The model is a discrete-iteration simulator: each iteration advances every in-batch sequence one decode step (or processes a prefill); the per-iteration time grows with batch size (more work per step → higher TPOT), which is what creates the latency-throughput dial.

1. What the output shows

== static vs continuous batching (Poisson load) ==
policy       thru(tok/s)  p50_TTFT  p99_TTFT  p50_TPOT completed
static               255     35774     75020        31       400
continuous           378        33      1669        31       400
  -> continuous: 1.5x throughput, 45x lower p99 TTFT

== goodput vs batch depth (SLO: TTFT<300ms, TPOT<40ms) ==
 max_batch  throughput   goodput
         4         209         3   <- batch too small: saturates -> TTFT SLO fails
         8         337        36   <- still saturating
        16         378       321       <- both SLOs met: goodput ~= throughput
        32         368         0   <- TPOT SLO blown (iter time 47ms > 40ms)
        64         350         0   <- deeper batch, worse goodput
goodput-optimal batch depth: 16

• Static suffers batch-formation delay (waits for a whole batch to arrive) and head-of-line blocking (the batch runs at its longest member's length while short requests waste their slots) → enormous TTFT. Continuous admits and retires per iteration → here 1.5× throughput and 45× lower p99 TTFT (WARMUP Ch. 5). At moderate load the headline win is latency; under saturation it becomes throughput.
• Goodput sweep — the headline result: raw throughput is highest at batch 32, but goodput collapses to zero there because per-token time (47 ms) blows the TPOT SLO. Too-small batches saturate and blow the TTFT SLO. Goodput peaks at an interior depth (16) — proof the objective is goodput, not max batch (WARMUP Ch. 6).

The per-iteration time scales with the configured max_batch (real engines graph-capture/pad decode kernels to the batch width), which is what makes a larger batch genuinely cost every user's TPOT. Production continuous-vs-static gains are often larger (3–10×) once variable arrivals and prefill interference — beyond this model — are included.

2. Reading order (solution.py)

1. Request / Poisson arrival generation.
2. simulate_static — fixed batches, run to completion.
3. simulate_continuous — the iteration loop: retire finished, admit waiting, step the batch; per-iteration time as a function of batch size.
4. The goodput accounting: a request's tokens count only if its TTFT and TPOT met the SLO.

3. Extensions

1. Chunked prefill: split long prefills across iterations so they don't stall decodes; measure the p99 TPOT improvement under a bimodal (short+long prompt) workload (WARMUP Ch. 5).
2. Prefix-cache integration: requests sharing a system prompt skip prefill (use Lab 01's idea) — watch TTFT collapse on cache hits.
3. SLO-aware admission: reject/queue requests when admitting them would blow the batch's TPOT SLO; show goodput protected under overload.
4. Prefill/decode disaggregation: separate prefill and decode "pools"; model the KV hand-off cost; compare interference vs the chunked-prefill approach.

4. Common pitfalls

1. Constant per-iteration time regardless of batch — then there's no latency-throughput dial and goodput never turns over. Per-iteration time must grow with batch size (the simulator models this).
2. Counting throughput as goodput — goodput must filter by SLO compliance, or the whole lesson disappears.
3. Poisson vs back-to-back arrivals — back-to-back inflates throughput unrealistically (WARMUP Ch. 6); the sim uses Poisson.
4. Ignoring prefill cost — a long prefill should visibly stall the iteration it enters (motivating chunked prefill).

5. What this lab proves about you

You can quantify the two serving truths that matter most to the P&L: continuous batching's throughput win and the goodput-optimal operating point. You can walk an exec through "why we don't just crank the batch size" with a graph, and you can design admission/SLO policy from the goodput curve — the platform-leadership deliverable this JD hires for.

Phase 08 — Distributed Systems & HPC Collectives

Difficulty: ⭐⭐⭐⭐⭐ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], Distributed Systems Engineer, HPC/ML-Systems Engineer Hardware needed: none — the collectives lab runs over real OS sockets on one machine (multi-process); the topology lab is a simulator

Why This Phase Exists

The JD wants "distributed systems, HPC" and "AI infrastructure at production scale." A single GPU is never enough at the frontier: training a large model and serving it with tensor parallelism both turn many GPUs into one logical device, and the glue is collective communication — all-reduce, all-gather, reduce-scatter — running over NVLink and InfiniBand. When training "mysteriously" runs at half speed, or a node failure hangs an 1,024-GPU job, the cause is almost always in this layer.

You'll implement ring all-reduce over real sockets (the algorithm NCCL uses) and benchmark it against the theoretical bandwidth bound, then build a topology-aware collective simulator that shows why placement (Phase 06) and bandwidth (Phase 01 Ch. 9) decide distributed throughput — and where failures turn into hangs.

Concepts

• Why distribute: model/data don't fit one GPU; the parallelism taxonomy
  • Data parallel (replicate model, split batch, all-reduce gradients)
  • Tensor parallel (split each layer across GPUs, all-reduce per layer — NVLink-bound)
  • Pipeline parallel (split layers into stages, micro-batch, bubble overhead)
  • Expert/sequence/context parallel (overview) and how they compose (3D/4D parallelism)
• Collective operations: broadcast, reduce, all-reduce, all-gather, reduce-scatter, all-to-all — and their communication volumes
• Ring all-reduce: the algorithm and its 2(N-1)/N × data bandwidth optimality; reduce-scatter + all-gather decomposition
• Tree vs ring vs hierarchical collectives; latency vs bandwidth regimes; small vs large messages
• NCCL: what it does, how it picks algorithms/topology, NCCL_DEBUG, the env knobs that matter
• Topology: NVLink/NVSwitch islands, PCIe, InfiniBand/RoCE, rails, GPUDirect RDMA; why placement (Phase 06) sets collective bandwidth
• Overlap: computation/communication overlap, bucketing gradients, the critical path
• Failure domains at scale: a single GPU/NIC failure hangs a synchronous collective; detection, checkpointing, elastic/fault-tolerant training; the "1% of nodes fail daily at 10k-GPU scale" reality
• Serving-side distribution: tensor-parallel inference, KV-cache sharding, prefill/decode disaggregation traffic (Phase 07)

Labs

Lab 01 — Ring All-Reduce Over Real Sockets (Python)

Lab 02 — Topology-Aware Collective Simulator (Python)

Deliverables Checklist

• Ring all-reduce: correct across N processes; bandwidth report matches the bound
• You can derive the 2(N-1)/N all-reduce volume and explain why it's N-independent
• Topology sim: placement-comparison table reproduced and explained
• You can explain TP/DP/PP and why TP must stay on NVLink
• You can describe how a single-GPU failure hangs a synchronous job and the mitigations
• One page: "our platform's multi-GPU/multi-node strategy" (feeds Phase 12)

Interview Relevance

• "Derive ring all-reduce's communication cost. Why is it bandwidth-optimal?"
• "Compare data, tensor, and pipeline parallelism. Which is NVLink-bound and why?"
• "Your 512-GPU training run is at 60% MFU. Where do you look?" (overlap, topology, stragglers)
• "One GPU fails in a 1024-GPU job. What happens, and how do you make it survivable?"
• "Why does tensor-parallel placement across nodes destroy throughput?"

Warmup Guide — Distributed Systems & HPC Collectives

Zero-to-expert primer for Phase 08. Builds on Phase 01 Ch. 9 (interconnect) and Phase 06 (placement). No distributed-training background assumed. By the end you can derive collective costs, reason about parallelism strategies, and design for failure at scale.

Table of Contents

• Chapter 1: Why Distribute, and the Parallelism Taxonomy
• Chapter 2: Collective Operations — the Vocabulary
• Chapter 3: Ring All-Reduce — the Bandwidth-Optimal Algorithm
• Chapter 4: Tree, Hierarchical, and When Each Wins
• Chapter 5: NCCL — the Production Collective Library
• Chapter 6: Topology — NVLink Islands, InfiniBand, and Placement
• Chapter 7: Overlap — Hiding Communication Behind Compute
• Chapter 8: Failure Domains at Scale
• Chapter 9: Distribution on the Serving Side
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: Why Distribute, and the Parallelism Taxonomy

Zero background: one GPU has finite HBM (80 GB) and finite compute. A 175B-parameter model in FP16 is 350 GB of weights alone — it doesn't fit. Even when a model fits, one GPU's throughput caps training time at "months." So we split the work across many GPUs and make them cooperate. The cooperation is the hard part, and it's all communication.

The four ways to split (you must be able to contrast these cold):

• Data parallel (DP): every GPU holds a full copy of the model, processes a different slice of the batch, and they all-reduce gradients each step to stay in sync. Simple, the workhorse. Communication: one all-reduce of the full gradient per step. Scales batch size, not model size.
• Tensor parallel (TP): split each layer's matrices across GPUs (e.g., each GPU computes part of every attention head), with an all-reduce inside every layer to combine partial results. Lets a model bigger than one GPU run — but the per-layer all-reduce is on the critical path, so TP must live on NVLink (Ch. 6); across slow links it's catastrophic.
• Pipeline parallel (PP): split the model's layers into stages on different GPUs; micro-batches flow through like an assembly line. Communication is just activations between adjacent stages (cheap), but you pay a pipeline bubble (stages idle at fill/drain). Lets very deep models run.
• Expert / sequence / context parallel: (overview) shard along other axes — MoE experts, the sequence dimension for long context. They compose: frontier training is "3D/4D parallelism" — e.g., TP within a node, PP across a few nodes, DP across the rest.

The mental model: DP trades communication for batch scale; TP trades communication for model scale (on fast links only); PP trades latency (bubble) for model depth. Choosing and composing them is a topology problem (Ch. 6).

Chapter 2: Collective Operations — the Vocabulary

A collective is a communication pattern over a group of ranks. The ones that matter, with their communication volume (V = data size per rank, N = ranks):

The headline fact: all-reduce = reduce-scatter + all-gather, and each of those moves (N-1)/N · V, so all-reduce moves ≈ 2V per rank, independent of N — the result you'll prove in Lab 01. This is why data-parallel training scales: adding GPUs doesn't increase per-GPU communication volume (only latency terms grow). All-reduce of gradients is the single most important collective in ML.

Chapter 3: Ring All-Reduce — the Bandwidth-Optimal Algorithm

The algorithm NCCL uses for large messages, and the centerpiece of Lab 01.

Setup: N ranks in a logical ring (each talks to its left and right neighbor). The gradient vector is chopped into N chunks. Two phases:

Phase 1 — reduce-scatter (N-1 steps): in step k, each rank sends one chunk to its right neighbor and receives a chunk from its left, adding the received chunk to its local copy. After N-1 steps, each rank holds the fully reduced sum of one chunk (a different chunk per rank).

Phase 2 — all-gather (N-1 steps): same ring motion, but now ranks forward the completed chunks (no addition) until everyone has all N reduced chunks.

The cost (the result to know cold): each rank sends/receives (N-1) chunks of size V/N in each phase = 2 phases × (N-1) × V/N = 2·(N-1)/N · V ≈ 2V per rank. As N→∞ this approaches exactly 2V, independent of N. Compare the naive "all-gather everything then sum locally": N·V per rank — N times worse. That N-independence is why ring all-reduce made large-scale data-parallel training practical (Baidu's 2017 result that propagated into every framework).

Why bandwidth-optimal: information-theoretically, every rank's contribution must reach every other rank, and each rank must receive the full reduced result — 2V is the lower bound on bytes that must cross each rank's links. Ring hits it. (The cost is N-1 latency hops per phase, so ring is bandwidth-optimal but latency-linear — hence trees for small messages, Ch. 4.)

Chapter 4: Tree, Hierarchical, and When Each Wins

Ring is optimal in bandwidth but its latency grows linearly in N (2(N-1) sequential hops). For small messages (latency-dominated) or many ranks, other structures win:

• Tree all-reduce: reduce up a binary tree to the root, broadcast down. Latency O(log N) instead of O(N) — far better for small messages. Bandwidth is worse (root links are hot), but small messages don't saturate bandwidth anyway. NCCL auto-switches ring↔tree by message size.
• Hierarchical / two-level: reduce within each NVLink island first (fast), then all-reduce across islands over IB (one representative per node), then broadcast back down. This matches the bandwidth hierarchy (Ch. 6) — you do the big-volume reduction on the fast intra-node links and only the small cross-node step on the slow links. Lab 02 quantifies the win.
• Double binary tree (NCCL's default for many cases): two trees using complementary links, recovering bandwidth while keeping log-latency.

The takeaway for a leader: there's no single best collective algorithm — NCCL picks by message size, rank count, and topology, and your job is to give it a good topology (Ch. 6) and large enough messages (bucketing, Ch. 7), not to hand-pick algorithms.

Chapter 5: NCCL — the Production Collective Library

NVIDIA Collective Communications Library — the layer every framework (PyTorch DDP/FSDP, Megatron, DeepSpeed) calls for GPU collectives.

What it does that you won't reimplement: detects the topology (NVLink, NVSwitch, PCIe, IB rails) at init, builds rings/trees over it, picks algorithms by message size, uses GPUDirect RDMA (GPU↔NIC↔GPU without staging through host memory — the Phase 01 PCIe-cliff avoidance), overlaps with compute via CUDA streams, and handles multi-rail IB.

The operational knobs/signals you will use:

• NCCL_DEBUG=INFO — prints the topology and algorithm choices (first stop when collectives are slow).
• NCCL_ALGO, NCCL_PROTO — force ring/tree, LL/Simple protocols (debugging).
• NCCL_IB_*, NCCL_SOCKET_IFNAME, NCCL_P2P_* — interface/transport selection (the cause of half of "training is slow on new hardware" incidents — wrong NIC, P2P disabled, IB not used).
• NCCL_TIMEOUT / the framework's collective timeout — what turns a hang (Ch. 8) into a crash with a stack trace.

The leadership point: NCCL is a black box that's fast when the topology and config are right and silently 10× slow when they aren't. Knowing what it's doing (Ch. 3–4, 6) is what lets you debug "MFU is 50%, why?" instead of filing a ticket and waiting.

Chapter 6: Topology — NVLink Islands, InfiniBand, and Placement

This is where Phase 06 (placement) and this phase (collectives) fuse. The bandwidth hierarchy (Phase 01 Ch. 9), as the substrate for collectives:

within a GPU:    HBM           ~3.35 TB/s
GPU<->GPU (node): NVLink/NVSwitch ~900 GB/s   <- TP lives here
node<->node:     InfiniBand/RoCE ~50 GB/s    <- DP/PP cross this
host:            PCIe           ~64 GB/s

The consequences that decide real throughput:

• Tensor parallelism must stay within an NVLink island. TP does a per-layer all-reduce on the critical path; on NVLink that's ~microseconds, on IB it's ~20× slower and serializes every layer — a 2–5× total slowdown. This is why schedulers (Phase 06) gang-place TP groups on one NVLink-connected node, and why "my 8-GPU TP job is slow" is almost always "it got spread across two nodes."
• Data parallelism can cross nodes — its all-reduce is once per step and overlappable (Ch. 7), so IB bandwidth is tolerable; you still want hierarchical all-reduce (Ch. 4) to minimize cross-node volume.
• Rails and GPUDirect RDMA: multi-NIC nodes have "rails"; a topology-aware collective uses the NIC closest to each GPU (GPUDirect RDMA) to avoid PCIe hops. Misconfiguration here silently routes through host memory (the 50× cliff).

Composing parallelism to the topology is the frontier-training art: TP within node (NVLink), PP across a few nodes (cheap activation passing), DP across the rest (overlappable all-reduce). Lab 02 makes the placement-vs-throughput relationship quantitative.

Chapter 7: Overlap — Hiding Communication Behind Compute

Even optimal collectives cost time; the trick is to not wait for them.

• Gradient bucketing + overlap (DDP's core optimization): instead of one giant all-reduce after the whole backward pass, group gradients into buckets and all-reduce each bucket as soon as it's ready — overlapping communication of early-computed gradients with the still-running backward pass. Done right, communication is almost free (hidden behind compute). Bucket size is a tuning knob (too small = latency overhead per bucket; too large = less overlap).
• Computation/communication overlap generally: collectives run on a separate CUDA stream (Phase 05) so the compute stream proceeds; event edges enforce the dependency. The same stream-scheduling skills from Phase 05, at cluster scale.
• The critical-path mindset: at 60% MFU (model FLOPs utilization), the missing 40% is usually exposed communication (no overlap), pipeline bubbles, or stragglers (Ch. 8) — not raw kernel speed. Profiling distributed training is about finding what's not overlapped.

Chapter 8: Failure Domains at Scale

The chapter that separates "ran a 4-GPU job" from "operated a 10,000-GPU cluster." Synchronous collectives have a brutal property: every rank must participate, so any one failure hangs everyone.

• The hang: a GPU throws an Xid error (Phase 04/10), a NIC flaps, a process OOMs — that rank stops calling the collective, and all other ranks block forever inside NCCL waiting for it. No error, no crash — a hang. At 10k-GPU scale where ~1% of nodes have an issue daily, this happens constantly.
• Detection: collective timeouts (NCCL_TIMEOUT, framework watchdogs) turn hangs into crashes with a culprit rank; health monitoring (Phase 10) catches the failing GPU/NIC; "straggler" detection finds the slow-but-not-dead rank dragging the whole job to its speed.
• Recovery: checkpointing (the only real defense — periodic, async, sharded checkpoints so a failure costs minutes not days), elastic training (frameworks that can drop/re-add ranks and continue), and redundancy/hot spares (replace the failed node, restore from checkpoint, resume). The checkpoint interval is an explicit cost/risk tradeoff you own: more frequent = less lost work per failure, more overhead.
• Blast radius: placement (Ch. 6) determines how much one failure costs — a failure in a TP group stalls that whole group; good placement minimizes the blast radius. Lab 02's extension models this.

The leadership reframe: at scale, failures are not exceptional, they're the steady state. You design the training/serving platform assuming continuous partial failure — checkpoint cadence, fast detection, elastic recovery, and spare capacity are first-class features, not afterthoughts. This is also the sovereign/on-prem story (Phase 07 Ch. 9): air-gapped clusters can't lean on cloud auto-healing.

Chapter 9: Distribution on the Serving Side

Collectives aren't just for training. Large-model serving (Phase 07) is distributed too:

• Tensor-parallel inference: a model too big for one GPU is TP-sharded across an NVLink island; every forward pass does per-layer all-reduces — so serving a 70B+ model has the same NVLink-island requirement as TP training.
• KV-cache sharding: the KV cache (Phase 07) is split across the TP group; attention does its own communication.
• Prefill/decode disaggregation (Phase 07 Ch. 5): separate GPU pools for prefill and decode means shipping the KV cache between them over the interconnect — a new, large communication pattern that topology (Ch. 6) must accommodate.
• Expert parallelism for MoE serving: all-to-all to route tokens to experts — the all-to-all collective at inference time, latency-sensitive.

So the topology and collective skills here apply directly to the JD's serving mandate: a platform serving large models is a distributed system, and its throughput is gated by the same NVLink/IB hierarchy and the same overlap discipline.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02 (implement the algorithm, then reason about placing it on a topology).

Lab 01 (ring all-reduce over sockets):

1. python solution.py --world 4; confirm correctness, then read the reduce-scatter/all-gather phases against Ch. 3.
2. Verify the byte count: each rank moves ≈2(N-1)/N · V — measure it, don't trust the formula. Compare to the naive all-gather-sum (N·V).
3. Extensions: tree all-reduce (measure latency on small vs large messages — Ch. 4's crossover); bucketing/overlap; then the failure demo — kill a rank, watch the hang, add a timeout (Ch. 8).

Lab 02 (topology simulator):

1. python solution.py; read the bandwidth-hierarchy cost model against Ch. 6.
2. Reproduce: TP within an NVLink island vs TP across nodes — the ratio is the NVLink/IB ratio, quantified. This is the Phase-06-meets-Phase-08 result.
3. Hierarchical all-reduce vs flat across nodes — Ch. 4's win, measured.
4. Extension: pipeline bubble model and failure-domain blast radius.

Success Criteria

• You can contrast DP/TP/PP and state which is NVLink-bound and why (Ch. 1)
• You can derive all-reduce = reduce-scatter + all-gather = 2V/rank and explain N-independence (Ch. 2–3) — and your Lab 01 byte count proves it
• You can say when tree beats ring and what hierarchical all-reduce buys (Ch. 4)
• You can name the NCCL signals/knobs for debugging a slow collective (Ch. 5)
• You can explain why TP across nodes is catastrophic, with the bandwidth ratio (Ch. 6) — Lab 02 quantifies it
• You can explain gradient overlap and diagnose low MFU (Ch. 7)
• You can describe how one failure hangs a job and the detection/recovery stack (Ch. 8)

Interview Q&A

Q1: Derive ring all-reduce's communication cost and explain why it's bandwidth-optimal. A: All-reduce = reduce-scatter + all-gather. Chop the V-byte vector into N chunks. Reduce-scatter: N-1 steps, each rank sends/receives one V/N chunk per step and accumulates → (N-1)·V/N moved, ending with each rank owning one fully- reduced chunk. All-gather: another N-1 steps of V/N chunks → another (N-1)·V/N. Total ≈ 2·(N-1)/N · V ≈ 2V per rank, independent of N. It's bandwidth-optimal because every rank must export its contribution and import the full reduced result — 2V is the information-theoretic minimum bytes across each rank's links, and the ring saturates it. The catch: 2(N-1) sequential hops means latency grows with N, so for small messages a log-latency tree wins — which is why NCCL switches by message size.

Q2: Compare data, tensor, and pipeline parallelism — which is NVLink-bound? A: Data parallel replicates the model, splits the batch, and all-reduces gradients once per step — communication is overlappable with the backward pass and tolerates cross-node IB, so DP scales across nodes. Tensor parallel splits each layer's matrices across GPUs with an all-reduce inside every layer on the critical path — that frequency makes it NVLink-bound: on NVLink it's microseconds, across nodes on IB it's ~20× slower per layer and serializes the whole forward/backward, a 2–5× slowdown. Pipeline parallel splits layers into stages passing only activations between neighbors (cheap), trading a pipeline bubble (idle fill/drain) for depth. Frontier training composes them: TP within a node, PP across a few nodes, DP across the rest.

Q3: Your 512-GPU training run is at 60% MFU. Where do you look? A: The missing 40% is almost never raw kernel speed; it's exposed time. In order: (1) communication not overlapped — is gradient bucketing/overlap on? Profile for all-reduce on the critical path; (2) topology mistakes — is TP accidentally spanning nodes (NCCL_DEBUG=INFO shows the rings; check it's using NVLink/IB not PCIe/sockets), is GPUDirect RDMA active; (3) pipeline bubbles — micro-batch count too low for the PP depth; (4) stragglers — one slow GPU (thermal throttle, ECC retries) dragging every collective to its speed, found by per-rank step-time variance; (5) data loading / checkpoint stalls on the critical path. The discipline is "find what's not overlapped," via the distributed profiler and NCCL logs, not micro-optimizing kernels.

Q4: One GPU fails in a 1024-GPU synchronous job. What happens and how do you make it survivable? A: The failed rank stops entering the collective, so all 1023 others block forever inside NCCL — a silent hang, not a crash. Survivability is a stack: (1) detection — collective timeouts and watchdogs convert the hang into a crash naming the culprit rank; health monitoring (Xid/NIC) flags the failing hardware; straggler detection catches slow-not-dead ranks. (2) recovery — checkpointing is the real defense: frequent, async, sharded checkpoints so a failure costs minutes; elastic frameworks drop/re-add ranks; hot spares replace the node and resume from checkpoint. (3) blast-radius minimization — placement so one failure stalls the smallest group. At 10k-GPU scale ~1% of nodes fail daily, so this isn't exceptional handling — continuous partial failure is the design assumption, and checkpoint cadence is an explicit cost/risk lever I'd own.

Q5: Why does tensor-parallel placement across nodes destroy throughput, with numbers? A: TP all-reduces partial results inside every layer on the critical path. On NVLink (~900 GB/s, ~µs latency) that per-layer collective is negligible; across nodes on InfiniBand (~50 GB/s, ~µs but ~18× less bandwidth and extra hops) each all-reduce is ~20× slower and, because it's serial within each layer's forward and backward, it can't be hidden — so a model with dozens of layers pays that penalty dozens of times per step, often a 2–5× total slowdown. That's why the scheduler (Phase 06) must gang-place a TP group on one NVLink-connected node, and why "8-GPU TP job is slow" almost always means it got spread across two nodes — the single most common distributed-training misconfiguration.

Q6 (leadership): How do you architect a training platform to be efficient AND fault-tolerant at 10k-GPU scale? A: Two intertwined goals. Efficiency: compose parallelism to the topology (TP within NVLink islands, PP across a few nodes, DP across the rest), enforce it through topology-aware gang scheduling (Phase 06), maximize overlap (gradient bucketing, comm on separate streams), and monitor MFU with per-rank step-time variance to catch stragglers. Fault tolerance: assume continuous partial failure — async sharded checkpointing at a cadence tuned to the failure rate (minutes of lost work, not days), collective timeouts + health monitoring for fast detection with a named culprit, elastic training to continue through node loss, and hot-spare capacity for fast replacement. Both rest on observability (Phase 10): without per-rank metrics and NCCL visibility you can neither find the efficiency leaks nor detect the failures. And for sovereign/ on-prem (Phase 07/11), all of this must work without cloud auto-healing — the platform owns the recovery, not the cloud provider.

References

• Baidu, "Bringing HPC Techniques to Deep Learning" (ring all-reduce for DL, 2017) — https://andrew.gibiansky.com/blog/machine-learning/baidu-allreduce/
• Thakur, Rabenseifner, Gropp, "Optimization of Collective Communication Operations in MPICH" (2005) — the algorithm survey
• NVIDIA NCCL documentation + NCCL_DEBUG — https://docs.nvidia.com/deeplearning/nccl/
• Shoeybi et al., "Megatron-LM" (tensor parallelism) — https://arxiv.org/abs/1909.08053
• Narayanan et al., "Efficient Large-Scale Language Model Training on GPU Clusters" (3D parallelism / PTD-P) — https://arxiv.org/abs/2104.04473
• Rajbhandari et al., "ZeRO" (DeepSpeed memory-efficient DP) — https://arxiv.org/abs/1910.02054
• PyTorch DDP design note (gradient bucketing/overlap) — https://pytorch.org/docs/stable/notes/ddp.html
• Meta, "OPT-175B logbook" / "The Llama 3 Herd of Models" — real failure-rate accounts at scale
• Cross-track: Phase 01 WARMUP Ch. 9 (interconnect), Phase 06 WARMUP Ch. 7 (topology-aware placement)

🛸 Hitchhiker's Guide — Phase 08: Distributed Systems & HPC Collectives

Read this if: you can train on one GPU but "all-reduce," "tensor parallel," and "why is MFU 50%" are fog. Field notes; derivations in the WARMUP.

0. The 30-second mental model

Many GPUs become one via collectives (all-reduce gradients/activations) over a bandwidth hierarchy (NVLink ≫ InfiniBand ≫ PCIe). Ring all-reduce moves ~2V/rank regardless of N (bandwidth-optimal). Tensor parallel must live on NVLink; data parallel can cross nodes. At scale, any one failure hangs everyone — so checkpoint, detect, recover. Low MFU = exposed communication, bubbles, or stragglers, not slow kernels.

1. Parallelism card

Frontier = compose them: TP in node, PP across few nodes, DP across rest.

2. Collective cost card

all-reduce = reduce-scatter + all-gather = 2·(N-1)/N·V ≈ 2V per rank   (N-independent!)
naive all-gather-then-sum = N·V per rank                                (N× worse)
ring: bandwidth-optimal, latency O(N)   |   tree: bandwidth worse, latency O(log N)
NCCL auto-picks by message size; hierarchical = reduce in-node then across-node

3. Topology stakes (the numbers)

HBM 3.35 TB/s  >  NVLink 900 GB/s  >  PCIe 64 GB/s  ≈  InfiniBand 50 GB/s

TP across nodes instead of NVLink ≈ 2–5× total slowdown (per-layer all-reduce ×20 slower × dozens of layers, unhideable). The #1 distributed misconfig. Schedulers must gang-place TP groups on one NVLink island (Phase 06).

4. NCCL survival kit

NCCL is fast when topology+config are right, silently 10× slow when not. Know what it's doing.

5. Diagnosing low MFU (the flowchart)

MFU < 70%? ->
  comm not overlapped?   bucketing on? all-reduce on critical path? (most common)
  topology wrong?        TP spanning nodes? PCIe instead of NVLink/IB? GPUDirect off?
  pipeline bubbles?      too few micro-batches for PP depth
  stragglers?            per-rank step-time variance -> one slow GPU (thermal/ECC)
  data/checkpoint stall? on the critical path?

Find what's not overlapped. Kernel micro-opt is rarely the answer at scale.

6. Failure reality

At 10k GPUs, ~1% of nodes have an issue daily. A synchronous collective hangs on any failure (no crash — a hang). Defense stack:

• Detect: collective timeouts + watchdogs (hang→crash+culprit), health monitoring (Xid/NIC), straggler detection.
• Recover: async sharded checkpoints (cadence = cost/risk lever), elastic training, hot spares.
• Minimize blast radius: placement so one failure stalls the smallest group.

Failures are the steady state; design for continuous partial failure. (And for sovereign/on-prem, no cloud auto-heal — you own recovery.)

7. War stories

• "8-GPU TP job at half speed" — scheduler spread it across 2 nodes; per- layer all-reduce hit IB not NVLink. Topology constraint → fixed. The classic.
• "Training hung at 3am, no error" — one GPU Xid'd; 1023 ranks blocked in NCCL forever. No NCCL_TIMEOUT set. Added it; now it crashes with the culprit rank in seconds.
• "Slow only on the new cluster" — NCCL_SOCKET_IFNAME picked the mgmt NIC, not IB. One env var, 10× collective speedup.
• "MFU 55% and kernels look fine" — gradient overlap was off (custom training loop bypassed DDP's bucketing). Re-enabled; MFU → 78%.
• "One slow GPU tanked a 256-GPU job" — thermal-throttled card did ECC retries; every all-reduce waited for it. Straggler detection drained the node.

8. Exit bar

You can derive ring all-reduce cost, contrast DP/TP/PP and place each on the topology, debug a slow collective via NCCL signals, diagnose low MFU, and design for failure at scale. Plus you've built ring all-reduce over real sockets and a topology cost model. Next: Phase 09 — the keystone, where all of this hardware knowledge becomes a hardware-independent platform (the JD's core mandate).

Lab 01 — Ring All-Reduce Over Real Sockets (Python)

Phase: 08 — Distributed & HPC | Difficulty: ⭐⭐⭐⭐☆ | Time: 5–8 hours Language: Python (stdlib socket, multiprocessing) | Hardware: none (multi-process on one machine)

Concept primer: ../WARMUP.md Ch. 2–4.

Run

python solution.py --world 4          # 4 worker processes, ring all-reduce over TCP
python solution.py --world 8 --size 1000000

0. The mission

Implement the algorithm NCCL uses, over real OS sockets across real processes — so the communication is genuine, not simulated. Ring all-reduce in two phases (WARMUP Ch. 3):

1. Reduce-scatter (N-1 steps): ring-pass chunks, accumulating, until each rank owns one fully-reduced chunk.
2. All-gather (N-1 steps): ring-pass the completed chunks until everyone has all of them.

Then prove the cost: each rank moves ≈2(N-1)/N × V bytes — measured by counting actual socket bytes — and beats the naive all-gather-then-sum (N×V).

1. What the output shows

ring all-reduce: world=8, vector=80000 doubles
  correctness: PASS (every rank's result == sum 1..8 = 36)
  bytes/rank moved : ring=1,120,000   naive_allgather=4,480,000
  per-rank volume vs 2(N-1)/N*V bound: 1.00x  (1.00 = bandwidth-optimal)
  ring moves 4.0x fewer bytes than naive all-gather

• Correctness: every rank's vector equals the true element-wise sum across all ranks — the collective is exact.
• Bytes/rank: the ring moves ≈2(N-1)/N × V per rank, hitting the bandwidth-optimal bound. The key result — run --world 4 then --world 8: the ring's per-rank volume stays ≈2V (the factor barely changes from 1.5V to 1.75V) while naive all-gather grows as (N-1)·V, so the ring's advantage widens with N (2× at N=4, 4× at N=8). That N-independence is why ring all-reduce made large-scale data-parallel training practical (WARMUP Ch. 3).

2. Reading order (solution.py)

1. worker — each rank: connects to its ring neighbors over TCP, gets its data shard, runs the two phases.
2. reduce_scatter — the N-1 send/recv-and-add steps; note the chunk indexing (rank r sends chunk (r-step) % N).
3. all_gather — the N-1 forward-only steps.
4. The byte counters and the correctness check against a reference sum.

3. Extensions

1. Tree all-reduce: implement reduce-up/broadcast-down; benchmark latency vs ring on small messages (1 KB) and large (10 MB) — find the crossover (WARMUP Ch. 4).
2. Bucketing/overlap: split the vector into buckets, all-reduce each as it's "ready," overlapping with simulated backward compute; measure the hidden time (WARMUP Ch. 7).
3. Failure demo: have one rank sys.exit() mid-collective; watch the others hang in recv. Then add a socket timeout that turns the hang into a detected failure with the culprit rank (WARMUP Ch. 8).
4. Bandwidth-optimality check: vary world size 2/4/8/16; confirm per-rank volume stays ≈2V (N-independent) while a naive all-gather-sum grows as N·V.

4. Common pitfalls

1. Deadlock from ordered send/recv: if every rank sends-then-recvs to the same neighbor synchronously, you can deadlock on socket buffers. The lab uses a send thread (or non-blocking) — keep that structure.
2. Chunk index off-by-one: the rank-to-chunk mapping per step is the crux; get it wrong and the reduction is incorrect (the correctness assert catches it, but the index math is the lab's real content).
3. Float serialization: use a fixed format (struct/array) consistently across ranks, or you'll get garbage sums.
4. Counting bytes wrong: count payload bytes actually sent on the wire, not the logical vector size.

5. What this lab proves about you

You've implemented and measured the algorithm at the heart of every data-parallel training run, and you can derive its bandwidth-optimal cost from the byte counts you collected — not recite it. NCCL is no longer magic: it's this, with topology detection and GPUDirect RDMA on top.

Lab 02 — Topology-Aware Collective Simulator (Python)

Phase: 08 — Distributed & HPC | Difficulty: ⭐⭐⭐⭐☆ | Time: 4–6 hours Language: Python (stdlib) | Hardware: none

Concept primer: ../WARMUP.md Ch. 4, 6.

Run

python solution.py

0. The mission

Model a cluster's interconnect hierarchy (NVLink within a node, InfiniBand between nodes) and compute collective time for different placements — proving the Phase-06-meets-Phase-08 thesis: where you put a parallel group decides throughput more than which algorithm you use.

1. What the output shows

== all-reduce time for an 8-way tensor-parallel group (256 MB activation, x80 layers/step) ==
placement                   links          per-step (ms)   vs best
TP within 1 NVLink node     NVLink                  44.0      1.0x
TP split across 2 nodes     InfiniBand             753.9     17.1x  <- catastrophic
TP split across 4 nodes     InfiniBand             753.9     17.1x

== data-parallel gradient all-reduce, 4 nodes x 8 GPUs (256 MB) ==
flat ring over all 32 GPUs (IB)          : 10.5 ms
hierarchical (NVLink in-node, IB across)  : 1.6 ms   (6.7x faster)

== failure blast radius (1 GPU dies, synchronous job) ==
TP-within-node placement : 1 GPU failure stalls 8 GPUs (one TP group)
TP-across-nodes placement: 1 GPU failure stalls 32 GPUs (spans all nodes)

(The TP cost is multiplied by 80 layers because tensor parallelism all-reduces inside every layer — which is exactly why crossing nodes is catastrophic: you pay the slow-link penalty dozens of times per step, unhideable.)

• TP placement: tensor parallelism within one NVLink node is ~18× faster than spread across nodes — the WARMUP Ch. 6 result, quantified. Same algorithm, same data; only placement differs. This is why "my 8-GPU TP job is slow" is almost always "it got spread across nodes."
• Hierarchical all-reduce: reduce within NVLink islands first, then across nodes over IB — does the big-volume reduction on fast links, only the small cross-node step on slow links (WARMUP Ch. 4). ~6× faster than a flat ring.
• Blast radius: placement also decides how many GPUs one failure stalls (WARMUP Ch. 8).

2. Reading order (solution.py)

1. Link bandwidths (NVLink/IB/PCIe) and the allreduce_time cost model (2(N-1)/N · V / bandwidth + latency).
2. The placement scenarios — how the ranks map to nodes and which links the collective crosses.
3. hierarchical_allreduce_time — the two-level decomposition.
4. The blast-radius computation.

3. Extensions

1. Pipeline bubble model: compute PP efficiency = microbatches / (microbatches + stages - 1); show how more micro-batches shrink the bubble.
2. 3D parallelism placement: TP=8 (in node), PP=4 (across nodes), DP=rest; compute total step time and find the placement that minimizes it.
3. Rail/GPUDirect model: add a PCIe-staging penalty when GPUDirect RDMA is misconfigured; show the 50× host-bounce cliff.
4. Failure-domain optimization: given a job shape, find the placement that minimizes blast radius subject to the NVLink-island constraint.

4. Common pitfalls

1. Using bandwidth without latency — for small messages latency dominates and ring loses to tree (WARMUP Ch. 4); the model includes a per-hop latency term.
2. Forgetting that TP all-reduces per layer — the per-step TP cost is n_layers × per-layer-allreduce, which is what makes cross-node TP catastrophic (not a one-time cost).
3. Treating all GPUs as equidistant — the entire lesson is that they aren't.

5. What this lab proves about you

You can quantify the interconnect's effect on distributed throughput and explain why placement (Phase 06) and collectives (Phase 08) must be co-designed. When a training job underperforms, you reason from topology — "is the TP group on one NVLink island?" — before touching kernels. That's the systems-level judgment the JD's "HPC at production scale" requires.

Phase 09 — Hardware Abstraction & Portability (The Keystone Phase)

Difficulty: ⭐⭐⭐⭐⭐ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], Runtime/Platform Architect, Compiler/Backend Engineer Hardware needed: none — the HAL and its backends are real C, runnable on any machine

Why This Phase Exists

This is the JD's central mandate, verbatim: "a platform that operates independently of the underlying hardware" and "decouple software from hardware environments." Every other phase has been building the knowledge to do this credibly — you can't abstract hardware you don't understand (Phases 01–08), and you can't sell the abstraction without the security/licensing that makes it a product (Phase 11).

A hardware-abstraction layer (HAL) is the single most important architectural artifact a hardware-agnostic AI platform owns. Get it right and you can support NVIDIA, AMD, Intel, and a startup's accelerator behind one API, route each workload to the best backend, and capture the market gap that NVIDIA's lock-in creates. Get it wrong and you've built either a leaky abstraction that performs terribly or a lowest-common-denominator wrapper that delivers a fraction of each chip's capability.

You will build a real HAL in C — a stable API, a backend ABI, and dlopen-loaded vendor plugins (CPU-reference, a "fast" backend, a "limited" backend) — the literal architecture of decoupling software from hardware. Then you'll grapple with the hard parts: capability negotiation, the abstraction-vs-performance tradeoff, and versioning the ABI.

Concepts

• API vs ABI: the difference, why it's the crux of plugin systems, ABI stability rules (struct layout, versioning, symbol visibility)
• The HAL pattern: a stable front-end API + a backend interface (vtable of function pointers) + dynamic backend loading (dlopen/dlsym)
• Capability negotiation: backends declare what they support (dtypes, ops, features); the platform queries and routes — no lowest-common-denominator
• The portability landscape revisited (from Phase 03 Ch. 9): CUDA, ROCm/HIP, oneAPI/Level Zero, Metal, Vulkan/SPIR-V, and the compiler layer (Triton/MLIR) as a backend strategy
• The abstraction-vs-performance tension: leaky abstractions, escape hatches, per-backend fast paths; why "write once, run anywhere, fast" is the hard part
• Backend selection/routing: choosing a backend per workload by capability + benchmark (ties to Phase 01 Ch. 10, Phase 07 Ch. 8)
• Versioning a plugin ABI: how to evolve the interface without breaking shipped backends (reserved fields, version structs, capability flags)
• Conformance testing: one test suite, all backends must pass — the contract that makes the abstraction trustworthy
• Build/packaging: shipping a platform that loads the right backend for the customer's hardware (incl. air-gapped, Phase 07/11)
• The business case: why this layer is the platform's moat and how it's monetized (OEM integrations, the JD's commercialization mandate)

Labs

Lab 01 — A GPU HAL with dlopen Backends (C)

Lab 02 — Portability Strategy Decision Doc + Conformance Harness (Markdown + Python)

Deliverables Checklist

• HAL: builds, all backends load, conformance passes, capability routing works, ABI version checked
• You added a 4th backend without recompiling the core (proof the ABI is real)
• You can explain API vs ABI and the rules for evolving a plugin ABI
• ADR completed for a multi-vendor scenario with costed options
• Conformance harness runs and flags non-conformant backends
• One page: "our platform's hardware-abstraction architecture" (the Phase 12 capstone's backbone)

Interview Relevance

This phase is the JD's thesis. Expect the hardest, most senior questions here.

• "Design a hardware-abstraction layer for an AI platform. Walk me through the API/ABI boundary."
• "How do you avoid a lowest-common-denominator abstraction?"
• "How do you version a plugin ABI so a year-old vendor backend still loads?"
• "Build vs adopt: would you write your own compiler layer or use Triton/MLIR? Defend it."
• "How is decoupling software from hardware actually a defensible business?" (→ system-design/, Phase 11)

Warmup Guide — Hardware Abstraction & Portability (The Keystone)

Zero-to-expert primer for Phase 09 — the JD's central mandate. Builds on every prior phase (you abstract hardware you understand). By the end you can design, build, and version a hardware-abstraction layer and argue the portability strategy to a board.

Table of Contents

• Chapter 1: The Mandate — "Independent of the Underlying Hardware"
• Chapter 2: API vs ABI — the Crux of Every Plugin System
• Chapter 3: The HAL Pattern — Stable API + Backend Vtable + dlopen
• Chapter 4: Capability Negotiation — Beating the Lowest Common Denominator
• Chapter 5: The Abstraction-vs-Performance Tension
• Chapter 6: Versioning a Plugin ABI
• Chapter 7: Conformance Testing — What Makes an Abstraction Trustworthy
• Chapter 8: The Portability Landscape, as Backend Strategy
• Chapter 9: The Business Case — Why This Layer Is the Moat
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: The Mandate — "Independent of the Underlying Hardware"

The JD's first responsibility: "Define and execute the strategy to decouple software from hardware environments" and build "a platform that operates independently of the underlying hardware." This phase is the technical content of that sentence.

Why it's valuable: NVIDIA's dominance is a software moat (Phase 03 Ch. 10) — CUDA + 18 years of kernels. Customers and OEMs want out from under the lock-in: AMD MI300X has more HBM, Intel Gaudi is cheaper, startups promise better tokens/joule — but switching means rewriting the stack. A platform that lets software run unchanged across vendors captures exactly that pain. That's the commercial thesis, and a hardware-abstraction layer (HAL) is its core artifact.

The hard truth this phase teaches: a HAL is easy to build badly (a thin wrapper that runs everywhere at 20% of peak) and very hard to build well (one API, multiple backends, near-native performance, evolvable without breaking shipped vendor plugins). The difference is Chapters 2–7.

Chapter 2: API vs ABI — the Crux of Every Plugin System

The single most important distinction in this phase, and one most engineers blur.

• API (Application Programming Interface): the source-level contract — function names, signatures, types you compile against. "Call hal_matmul(a, b, c)." Breaking the API breaks compilation.
• ABI (Application Binary Interface): the binary contract — how those calls are actually made at the machine level: struct memory layout, field offsets, calling convention, symbol names, enum values, alignment. Breaking the ABI breaks already-compiled binaries — silently, often as corruption rather than a clean error.

Why this is the crux of a HAL: backends are separately compiled binaries (vendor .so files you dlopen at runtime — possibly built by a different team, a different compiler, a year apart). They share no source with your core; they share only the ABI. So the contract between your platform and a vendor backend is binary, and getting it stable is the whole game.

The ABI rules you must internalize (Lab 01 enforces them):

• Struct layout is the contract. If your core thinks Capabilities is {int version; int flags; int max_dtype;} and a backend was compiled when it was {int version; int flags;}, every field after flags is read at the wrong offset → garbage. Never reorder or insert fields in the middle; only append, and only with a version guard (Ch. 6).
• Pass a version field first, so the core can refuse incompatible backends before reading misaligned fields.
• Symbol visibility: a backend exports exactly one well-known entry point (e.g., hal_backend_register); everything else is hidden. The dlsym of that symbol is the only coupling.
• Use fixed-width types (int32_t, not int) and explicit enums with assigned values — int width and enum representation are platform-dependent.

Misconception: "if it compiles, the ABI is fine." No — ABI breaks happen without recompiling the other side. The backend compiled last year still dlopens fine and then reads your new struct wrong. ABI discipline is about the binaries you didn't rebuild.

Chapter 3: The HAL Pattern — Stable API + Backend Vtable + dlopen

The architecture (Lab 01 builds exactly this):

   application
        |  calls the stable API
        v
  +-----------+      hal_matmul(...), hal_alloc(...), hal_select_backend(...)
  |  HAL core |  <-- ONE stable API the app compiles against
  +-----------+
        |  dispatches through a vtable of function pointers
        v
  +------------------ backend ABI ------------------+
  | struct Backend {                                |
  |   uint32_t abi_version;                          |
  |   Capabilities caps;                             |
  |   void* (*alloc)(size_t);                        |  <-- function-pointer
  |   int   (*matmul)(const Tensor*, ...);           |      vtable: the ABI
  |   ...                                            |
  | };                                               |
  +-------------------------------------------------+
        ^            ^             ^
   cpu_ref.so    fast.so      limited.so      <-- dlopen'd at runtime

The three pieces:

1. Stable API — what the application calls. It never changes shape; that's the promise to application authors.
2. Backend interface (the vtable) — a struct of function pointers + a capability struct. This is the ABI the vendor plugins implement.
3. Dynamic loading — dlopen("backend_fast.so") then dlsym(h, "hal_backend_register") returns the filled vtable. The core never links against any backend at build time — which is what "decoupled from hardware" means mechanically: you can ship a new backend as a file, no core rebuild (Lab 01 extension 1 proves this).

This is precisely how real systems work: CUDA's driver API loads cubins; OpenCL/Level Zero load ICDs (Installable Client Drivers); LLVM loads target backends; even Postgres/Python load .so extensions. The pattern is universal; the discipline (Ch. 2, 6) is what's hard.

Chapter 4: Capability Negotiation — Beating the Lowest Common Denominator

The trap that kills naive HALs: if your API only exposes what every backend supports, you deliver the lowest common denominator — no FP8 because the old backend lacks it, no Tensor Cores because the CPU backend has none. Customers on good hardware pay for capabilities they can't use.

The fix is capability negotiation: each backend declares what it supports — dtypes (FP32/FP16/BF16/FP8/INT8), ops (matmul, attention, conv), features (unified memory, async, specific fusions), and performance hints — in a capability struct. The platform then:

• Queries capabilities at load time.
• Routes each workload to a backend that supports it (a FP8 request goes to an FP8-capable backend; a CPU-only deployment falls back to FP32).
• Exposes the union of capabilities to applications, with a clean failure ("no loaded backend supports FP8") rather than a silent downgrade.

This is the difference between a HAL that's a product and one that's a tax. Lab 01 implements capability routing: a request needing FP16 is steered away from the "limited" backend; the demo shows the router picking correctly. It also mirrors Phase 06's device-plugin capability advertising and Phase 07's multi-engine routing — same pattern, different layer.

Chapter 5: The Abstraction-vs-Performance Tension

The central design tension, and the honest hard part of "write once, run anywhere, fast."

• Too abstract → leaky or slow: a single generic matmul can't express each backend's fast path (Tensor Core fragment shapes, AMD's matrix cores, a startup's systolic array), so you get correct-but-mediocre performance.
• Too concrete → not portable: expose vendor-specific knobs and the application code becomes vendor-coupled again — you've lost the decoupling.

The resolutions real platforms use (you should be able to name all four):

1. Escape hatches: a generic path plus an optional backend_native_op() that lets performance-critical code call a backend's fast path when present, falling back otherwise. Pragmatic; slightly leaky.
2. High-level ops: abstract at the operation level (attention, layernorm), not the instruction level — so each backend implements the op optimally inside the abstraction (this is why serving engines expose "attention" not "matmul+softmax+matmul"). The most successful posture.
3. A compiler layer: don't hand-write per-backend kernels — lower a portable IR to each backend (Triton/MLIR, Phase 03 Ch. 9). Highest ceiling, highest cost.
4. Capability-gated fast paths: the abstraction picks the best implementation based on declared capabilities (Ch. 4).

The leadership judgment: abstract at the right altitude. Too low and you re-expose the hardware; too high and you can't reach peak. "Operation-level abstraction with capability-gated fast paths and an escape hatch" is the defensible default — and it's what the capstone (Phase 12) builds.

Chapter 6: Versioning a Plugin ABI

The chapter that separates a HAL that survives contact with vendors from one that breaks every release. You ship a backend ABI; vendors build .so files against it; then you need to evolve it without breaking their year-old binaries.

The techniques (Lab 01 extension 3 demonstrates them):

• A version field, first in the struct. The core reads abi_version before anything else and refuses (cleanly) backends it can't support — never reading misaligned fields from a too-new or too-old struct.
• Append-only struct growth. New capabilities/functions go at the end of the struct/vtable, behind a version check. An old backend (smaller struct) still has its early fields at the right offsets; the core just doesn't read the new tail for old abi_versions.
• Reserved fields. Ship uint32_t reserved[8] so future fields don't even change the struct size — old and new backends agree on layout.
• Capability flags over struct changes. Prefer adding a flag ("supports FP8") to an existing flags field over adding a struct field — flags are forward/backward compatible by construction.
• Semantic versioning of the ABI: major = breaking (old backends rejected), minor = additive (old backends still load). The core supports a range.

The real-world stakes: CUDA's forward-compatibility (Phase 03/04), OpenCL ICD versioning, the Linux kernel's stable-ABI-for-userspace promise — all are this chapter at industrial scale. A platform whose ABI breaks vendor backends every quarter has no vendor ecosystem, and the ecosystem is the product (Ch. 9).

Chapter 7: Conformance Testing — What Makes an Abstraction Trustworthy

An abstraction is only as good as its guarantee that every backend behaves the same. Without that, applications must special-case backends — and you've lost portability. The mechanism is a conformance suite: one test matrix every backend must pass.

• Correctness conformance: the same operations, same inputs, same expected outputs (within a stated tolerance — Phase 03's FMA/rounding caveat applies across vendors!) on every backend. Lab 01's demo runs one conformance test against all backends.
• Capability conformance: a backend that declares FP16 support must actually produce correct FP16 results — declarations are tested, not trusted.
• Performance conformance (the advanced tier, Lab 02 extension): a backend must hit X% of its roofline (Phase 01) on key ops, or it's "supported but not recommended" — preventing the "technically runs, uselessly slow" backend from poisoning the platform's reputation.

This is how Khronos (OpenCL/Vulkan conformance), the C/C++ standards (test suites), and every credible plugin ecosystem work: the conformance suite is the contract. As a leader you'd gate "certified backend" status on it — which is also a vendor-relations and commercial lever (Ch. 9).

Chapter 8: The Portability Landscape, as Backend Strategy

Revisiting Phase 03 Ch. 9 from the HAL builder's seat — these are your backend options:

• Vendor compute libraries directly (CUDA+cuBLAS/cuDNN, ROCm+rocBLAS, oneAPI+oneMKL, Metal Performance Shaders): each backend wraps the vendor's tuned kernels. Fastest to near-peak per vendor, most code (one backend per vendor), the pragmatic start.
• A portable kernel layer (Triton, increasingly multi-backend; or SYCL/oneAPI): write kernels once, lower to each backend. Less per-vendor code, inherits the compiler's backend coverage, ~90% of peak on supported ops (Phase 03 Ch. 9).
• A full compiler stack (MLIR dialects → each target): maximum control, maximum cost; the moonshot.
• SPIR-V/Vulkan compute: truly neutral, historically below peak for ML.

The platform reality (Phase 03 Ch. 10's three postures, now as architecture): blend — vendor-library backends for the hardware you must support today (time to market), a Triton/MLIR backend strategy for kernel velocity and future hardware, all behind your HAL (this phase) with capability routing (Ch. 4). The HAL is the constant; the backends are how you fill it. The AMD asymmetry from Phase 03 (no stable PTX-equivalent, per-gfx code objects) is a real backend-engineering cost your HAL must budget for.

Chapter 9: The Business Case — Why This Layer Is the Moat

Why a Head of Engineering — not just an architect — owns this:

1. It's the commercialization mandate. The JD lists "own platform adoption and commercialisation through partnerships" and "drive integrations with OEMs, cloud providers, and infrastructure partners." Every one of those integrations is a backend behind your HAL. The HAL's quality determines how fast you can onboard an OEM's chip — a direct revenue lever.
2. It inverts NVIDIA's moat. NVIDIA's lock-in is software; a platform that makes software hardware-agnostic sells freedom from that lock-in to every customer and every non-NVIDIA chipmaker who wants market access. Your HAL is the chipmakers' on-ramp to the software ecosystem they can't build alone.
3. Conformance is a commercial gate. "Certified backend" status (Ch. 7) is leverage in OEM negotiations and a quality promise to customers — you control who gets to be in the ecosystem and at what performance bar.
4. It's defensible. Anyone can wrap two SDKs in an afternoon; a performant, versioned, conformance-gated, multi-vendor HAL with an ecosystem of certified backends is years of accumulated engineering and relationships — the same compounding moat dynamic as CUDA, pointed the other way.

The capstone (Phase 12) composes this HAL with the scheduler (06), serving (07), and licensing (11) into the platform the JD describes. This phase is its spine.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02 (build the HAL, then write the strategy around it).

Lab 01 (the HAL in C):

1. make && ./hal_demo; read hal.h first — the ABI is the contract, study the Backend struct and Capabilities against Ch. 2–3.
2. Read hal.c (loader/router) then the three backends; trace one hal_matmul call from API → vtable → backend.
3. Reproduce: capability routing (FP16 request avoids the "limited" backend); ABI-version rejection of a stale backend.
4. Do extension 1 — add a 4th backend as a new .so without recompiling the core. When it loads and passes conformance, you've felt what "decoupled from hardware" means. Then extension 3 (ABI evolution) for Ch. 6.

Lab 02 (strategy ADR + conformance):

1. Read STRATEGY-ADR.md; it's a worked decision record — study its structure.
2. python conformance.py; understand the support matrix + pass/fail logic (Ch. 7).
3. Write your own ADR for the multi-vendor scenario — costed options, a phased rollout, the conformance gate. This is a Phase 12 leadership artifact in miniature.

Success Criteria

• You can state the API/ABI difference and the rules that keep an ABI stable (Ch. 2)
• You can draw the HAL pattern (API + vtable + dlopen) and explain why the core never links a backend (Ch. 3)
• You can explain capability negotiation and why it beats LCD (Ch. 4)
• You can name the four resolutions to the abstraction-vs-performance tension and pick the default (Ch. 5)
• You can version a plugin ABI so a year-old backend still loads (Ch. 6) — and your Lab 01 extension proves it
• You can explain why conformance testing makes the abstraction trustworthy (Ch. 7)
• You can argue the portability-posture blend and the business case (Ch. 8–9)

Interview Q&A

Q1: Design a hardware-abstraction layer for an AI platform. Walk me through the API/ABI boundary. A: Three pieces. A stable API the application compiles against (hal_alloc, hal_matmul, hal_select_backend) — its shape never changes, that's the promise to app authors. A backend ABI — a struct of function pointers (the vtable) plus a capability struct — that vendor plugins implement; this is the binary contract, so it's governed by ABI rules: version field first, fixed-width types, append-only struct growth, reserved fields, one exported registration symbol. And dynamic loading — the core dlopens each backend .so and dlsyms its registration function to get the filled vtable; the core links no backend at build time, which is the mechanical meaning of "decoupled from hardware" — a new chip is a new file. The API/ABI split is the crux because backends are separately compiled binaries sharing only the ABI, so a year-old vendor .so must still load correctly — which drives every versioning decision.

Q2: How do you avoid a lowest-common-denominator abstraction? A: Capability negotiation. Each backend declares what it supports — dtypes, ops, features, perf hints — in a capability struct read at load time. The platform routes each workload to a backend that supports it (FP8 work to an FP8-capable backend), exposes the union of capabilities to applications, and fails cleanly ("no backend supports FP8") rather than silently downgrading. Combined with operation-level abstraction (expose attention, not matmul+softmax) so each backend implements ops with its fast path inside the abstraction, plus optional native escape hatches for the last mile — you get portability without forcing everyone down to the weakest backend's feature set. The anti-pattern is an API that only exposes the intersection of all backends; that makes good hardware pay for the worst backend's limits.

Q3: How do you version a plugin ABI so a year-old vendor backend still loads? A: Put an abi_version field first and check it before reading anything else, so an incompatible backend is rejected cleanly instead of read at wrong offsets. Grow the interface append-only — new functions/fields at the end of the vtable/struct, gated by version, so an old (smaller) backend's early fields stay at correct offsets and the core simply doesn't read the new tail for old versions. Ship reserved fields so additions don't even change struct size. Prefer capability flags over new struct fields (flags are compatible by construction). Use semantic versioning — major = breaking (old backends rejected), minor = additive (old still load) — and have the core accept a range. This is exactly how OpenCL ICDs, CUDA forward-compat, and the Linux userspace-ABI promise work; breaking vendor backends every release means no vendor ecosystem, and the ecosystem is the product.

Q4: Build vs adopt — would you write your own compiler layer or use Triton/MLIR for kernels? A: Default to adopt. MLIR exists precisely so platforms don't hand-roll an IR: dialects, pass infrastructure, and existing lowerings to NVPTX/ROCm are free; Triton gives ~90% of hand-tuned on the kernels LLMs need with Python-level authoring and multi-backend lowering. Writing a full compiler re-pays all of that for marginal fit gains and an unbounded maintenance bill, and it slows your ability to absorb the next FlashAttention-shaped innovation (ecosystem velocity). I'd build my own only for a genuinely novel execution model the dialects can't express. The pragmatic architecture is a blend: vendor-library backends for hardware I must support today (time to market), a Triton/MLIR backend for kernel velocity and future chips, all behind my own HAL with capability routing — the HAL is mine (it's the moat), the kernel compiler I adopt.

Q5: How is "decoupling software from hardware" actually a defensible business? A: NVIDIA's dominance is a software moat — CUDA plus two decades of kernels — so customers and non-NVIDIA chipmakers share a pain: switching hardware means rewriting the stack. A platform that runs software unchanged across vendors sells freedom from that lock-in to customers and market access to every AMD/Intel/ startup chip that can't build a CUDA-class ecosystem alone — they become backends behind the HAL. The defensibility compounds: a performant, versioned, conformance-gated, multi-vendor HAL with an ecosystem of certified backends and OEM relationships is years of engineering and trust, not an afternoon's wrapper. Conformance certification is a commercial gate you control; OEM integrations are revenue; and the whole thing is the inverse of CUDA's moat — equally compounding, pointed at the incumbent. That's why the JD frames it as both an engineering and a commercialization mandate, and why the Head of Engineering owns it.

References

• Ulrich Drepper, "How To Write Shared Libraries" (the ABI/dlopen bible) — https://www.akkadia.org/drepper/dsohowto.pdf
• man dlopen, man dlsym — the loader API you'll build on
• Khronos OpenCL ICD + conformance model — https://www.khronos.org/conformance/
• Linux kernel "stable API nonsense" (Greg KH) — the opposite philosophy, instructive on why userspace ABI stability matters — https://www.kernel.org/doc/html/latest/process/stable-api-nonsense.html
• Intel oneAPI / Level Zero spec (a real cross-vendor HAL) — https://spec.oneapi.io/level-zero/latest/
• LLVM "Writing an LLVM Backend" (the compiler-as-backend posture) — https://llvm.org/docs/WritingAnLLVMBackend.html
• Triton + MLIR (Phase 03 references) — https://triton-lang.org , https://mlir.llvm.org
• Architecture Decision Records (Michael Nygard) — https://cognitect.com/blog/2011/11/15/documenting-architecture-decisions
• Cross-track: Phase 03 WARMUP Ch. 9–10 (the compiler layer), Phase 12 (the capstone this spine supports)

🛸 Hitchhiker's Guide — Phase 09: Hardware Abstraction (The Keystone)

Read this if: you've heard "decouple software from hardware" and want to know what it mechanically means. This is the JD's thesis. Field notes; derivations in the WARMUP.

0. The 30-second mental model

A HAL = stable API (what apps call) + backend ABI (a vtable of function pointers + capability struct that vendor plugins implement) + dlopen (load backends at runtime; core links none of them). Beat lowest-common-denominator with capability negotiation. Keep the ABI loadable across years with version-first, append-only, reserved fields. Make it trustworthy with a conformance suite. This layer is the platform's moat (Ch. 9).

1. API vs ABI (the thing everyone blurs)

• API = source contract (names, signatures). Breaking it breaks compiles.
• ABI = binary contract (struct layout, offsets, calling convention, symbol names, enum values). Breaking it breaks already-built binaries — silently.
• Backends are separately-compiled .so files sharing only the ABI with your core. That's why ABI discipline is the whole game.

ABI rules: version field FIRST · fixed-width types (int32_t) · append-only struct growth · reserved[N] fields · one exported symbol · enums with explicit values.

2. The pattern (Lab 01 in one diagram)

app -> stable API -> HAL core -> vtable dispatch -> backend.so (dlopen'd)
                                  struct Backend { abi_version; caps; (*matmul); ... }
cpu_ref.so   fast.so   limited.so   <- swap/add WITHOUT recompiling the core

Same pattern as CUDA driver/cubins, OpenCL ICDs, LLVM backends, Postgres extensions. Universal mechanism; the discipline is what's rare.

3. Capability negotiation (beat the LCD)

Each backend declares dtypes/ops/features. Platform queries → routes work to a capable backend → exposes the union, fails cleanly on the unsupported. Don't expose only the intersection of all backends — that taxes good hardware for the worst backend's limits. (Same shape as Phase 06 device-plugin advertising and Phase 07 engine routing.)

4. Abstraction vs performance (the hard part)

Too abstract = slow/leaky; too concrete = vendor-coupled. Four fixes:

1. Escape hatches (generic path + optional native fast path)
2. Operation-level ops (attention, not matmul+softmax) ← the winner
3. Compiler layer (Triton/MLIR lowering) ← highest ceiling, highest cost
4. Capability-gated fast paths

Default: operation-level + capability-gated fast paths + escape hatch. Abstract at the right altitude.

5. Versioning so vendors don't hate you

version-first check · append-only growth · reserved fields · flags over fields · semver (major=breaking, minor=additive), core accepts a range. Break vendor backends every quarter → no ecosystem → no product. The ecosystem IS the product.

6. Conformance (what makes it trustworthy)

One test matrix, all backends pass: correctness (with cross-vendor rounding tolerance — Phase 03), declared-capability verification, and (advanced) a performance tier (X% of roofline or "supported, not recommended"). "Certified backend" status is both a quality promise and an OEM-negotiation lever.

7. Backend strategy (Phase 03 Ch. 9 from the builder's seat)

Blend: vendor backends now + Triton/MLIR for velocity, all behind YOUR HAL. Budget for the AMD asymmetry (no PTX-equivalent; per-gfx rebuilds — Phase 03).

8. The business case (why a Head of Eng owns this)

• OEM integrations = backends behind your HAL = direct revenue lever (the JD's commercialization mandate).
• It inverts NVIDIA's software moat — sells lock-in freedom to customers and market access to non-NVIDIA chips.
• Conformance certification = a commercial gate you control.
• A performant, versioned, conformance-gated multi-vendor HAL is years of moat, not an afternoon's wrapper.

9. War stories

• "Our HAL ran everywhere at 25% of peak" — abstracted at the instruction level (generic matmul), couldn't reach Tensor Cores. Re-abstracted at op level (attention) with capability-gated fast paths. Peak recovered.
• "A vendor's old .so corrupted memory after our release" — we inserted a struct field in the middle; their year-old binary read every later field at the wrong offset. Now: append-only + reserved fields + version gate. Never again.
• "Backend declared FP16, produced garbage" — declarations were trusted, not tested. Conformance suite now verifies every declared capability.
• "We shipped one engine, lost a deal needing AMD" — no HAL, NVIDIA-coupled. The HAL is exactly the thing that would have made that deal a config change.

10. Exit bar

You can design and build a versioned, dlopen-based HAL with capability routing and conformance; explain API vs ABI and ABI evolution cold; pick the portability-posture blend; and argue the moat. You've added a backend without recompiling the core. This is the JD's thesis, in your hands. Next: Phase 10 — making all of it observable and operable in production.

Lab 01 — A GPU HAL with dlopen Backends (C)

Phase: 09 — Hardware Abstraction (Keystone) | Difficulty: ⭐⭐⭐⭐⭐ | Time: 6–10 hours Language: C11 + POSIX dlopen | Hardware: none

Concept primer: ../WARMUP.md Ch. 2–7.

Run

make            # builds the HAL core, 3 backend .so files, and the demo
./hal_demo

0. The mission

Build the literal architecture of "decouple software from hardware" (WARMUP Ch. 1): a stable API, a versioned backend ABI, and three dlopen-loaded vendor backends. The core links none of the backends — it loads them at runtime, queries their capabilities, and routes work to a capable one.

1. What the output shows

loaded backend: cpu_ref   (abi v2, dtypes=F32, perf=1)
loaded backend: fast      (abi v2, dtypes=F32|F16, perf=10)
loaded backend: limited   (abi v2, dtypes=F32, perf=2)
rejected backend: stale   (abi v1 != core v2) — ABI version guard works

== conformance: matmul FP32 vs reference, all backends ==
  cpu_ref : PASS
  fast    : PASS
  limited : PASS

== capability routing ==
  FP32 matmul -> fast      (highest perf among F32-capable)
  FP16 matmul -> fast      (only F16-capable backend)
  FP8  matmul -> (none)    cleanly refused: no loaded backend supports FP8

• Three backends loaded from .so files the core never linked against — that's the decoupling, mechanically (WARMUP Ch. 3).
• ABI version guard rejects a backend built against an incompatible ABI (WARMUP Ch. 2, 6) — the discipline that lets a year-old vendor binary still load (or be cleanly refused).
• Conformance: one test, all backends pass (WARMUP Ch. 7).
• Capability routing: FP16 work goes only to the FP16-capable backend; FP8 is refused cleanly instead of silently downgraded (WARMUP Ch. 4 — beating the lowest common denominator).