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