🛸 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.