The GPU Hardware Model

You are not writing “GPU code.” You are writing code for a very specific machine with a very specific dance: 132 (or 148) small computers, each with its own scratchpad, all sharing a fat but slow global memory, all fed by a hierarchy of async engines. Get the picture right and every optimization will feel obvious.

1. The one-slide anatomy

┌─────────────────────────────────────────────────────────────┐
│ GPU (H100 SXM5 example)                                     │
│                                                             │
│   HBM3 : 80 GB @ 3.35 TB/s                                  │
│                       │                                     │
│                    L2 cache : ~50 MB (shared across SMs)    │
│                       │                                     │
│   ┌───────────────────┴──────────────────────────────┐      │
│   │ 132 × Streaming Multiprocessor (SM):             │      │
│   │                                                  │      │
│   │   • 4 warp schedulers, 32 threads / warp         │      │
│   │   • 65,536 × 32-bit registers (256 KB)           │      │
│   │   • Shared memory + L1 : 228 KB configurable     │      │
│   │   • 4 × 4th-gen Tensor Cores (per SM)            │      │
│   │   • TMA engine (Hopper+), DMA async copies       │      │
│   └──────────────────────────────────────────────────┘      │
└─────────────────────────────────────────────────────────────┘\n```\n\nSame diagram, Blackwell (B200) delta:\n- 148 SMs, 5th-gen tensor cores\n- **TMEM: 256 KB dedicated Tensor Memory per SM** (in addition to SMEM)\n- HBM3e ~8 TB/s\n- `tcgen05` PTX instructions, 2-CTA MMA mode, hardware decompression engine\n- Native FP4 (NVFP4 & MXFP4)\n\n## 2. The execution units — words you must own\n\n- **Thread**: the smallest schedulable unit. Owns registers.\n- **Warp**: 32 threads executing in lockstep (SIMT). The real scheduling unit. Divergence within a warp = serialization.\n- **Warp scheduler**: picks a ready warp each cycle. 4 per SM on H100/B200. Occupancy exists to give schedulers something to switch to when a warp stalls on memory.\n- **Warp group** (Hopper+): 4 warps (128 threads) that cooperate on async tensor-core ops (`wgmma`). This is the granularity FA3 optimizes around.\n- **Block / CTA** (Cooperative Thread Array): threads that share an SM and can synchronize + share SMEM. Up to 1024 threads.\n- **Thread-block cluster** (Hopper+): a set of blocks pinned across nearby SMs that can access each other's SMEM via **DSMEM (distributed shared memory)**. Enables cross-SM cooperation without going through global memory.\n- **Grid**: the whole launch. Blocks in a grid do **not** synchronize (except via cooperative launch or global memory + atomics).\n\n## 3. The memory hierarchy — latencies + bandwidths\n\nInternalize this table. It is the source of every optimization.\n\n| Level | Size (H100) | Bandwidth | Latency (approx) | Scope |\n|---|---|---|---|---|\n| Register file | 256 KB/SM | ~ TB/s / SM | 1 cycle | Per thread |\n| SMEM / L1 | 228 KB/SM | ~ 20 TB/s aggregate | ~30 cycles | Per block |\n| TMEM (Blackwell) | 256 KB/SM | 16 TB/s R / 8 TB/s W | ~SMEM-ish | Per block (tensor-core operand path) |\n| L2 | ~50 MB | ~5 TB/s | ~200 cycles | GPU-wide |\n| HBM | 80–192 GB | 3.35–8 TB/s | ~400–800 cycles | GPU-wide |\n| PCIe / NVLink | — | 64 GB/s (Gen5) / 900 GB/s (NVL) | ~µs | Off-GPU |\n\n**Consequence:** any workload with intensity below the ridge is a memory-hierarchy problem, not a compute problem. Fusion exists because the round-trip HBM → SMEM → tensor core → SMEM → HBM is where time actually goes.\n\n## 4. Tensor cores — the only compute worth optimizing\n\nRule of thumb: on modern GPUs, tensor cores deliver **~ 10–16× more BF16/FP16 FLOPs** than CUDA cores. If your inner loop is not `wmma` / `wgmma` / `tcgen05`, you are running on a fraction of the chip.\n\n| Gen | Instructions | Sync | Precision highlights |\n|---|---|---|---|\n| Volta | `mma.sync` (m8n8k4) | sync | FP16 |\n| Ampere | `mma.sync` (m16n8k16) | sync | BF16, TF32, INT8 |\n| Hopper | **`wgmma`** (warp-group, async) | async | FP8 (E4M3/E5M2), BF16 |\n| Blackwell | **`tcgen05`**, 2-CTA MMA | async | **FP4 (NVFP4/MXFP4)**, FP6, FP8 |\n\nThe move from sync (`mma.sync`) to async (`wgmma`/`tcgen05`) is *the* Hopper/Blackwell story: while the tensor core computes tile *k*, TMA is DMA'ing tile *k+1* into SMEM and softmax is running on the previous output. Overlap is the game.\n\n## 5. Async engines you must name\n\n- **TMA (Tensor Memory Accelerator, Hopper+)**: a hardware engine that copies multi-dimensional tiles between GMEM and SMEM. One thread issues, engine handles addressing + coalescing + strides. Frees warps to compute.\n- **`cp.async` (Ampere)**: async global→shared copy that predates TMA. Still ubiquitous in Ampere-targeted CUTLASS.\n- **Async barrier / mbarrier**: hardware barrier (`arrive/wait`) used to signal when TMA/`wgmma` complete. Read: *asynchronous execution model* section of the PTX ISA.\n- **DSMEM (Hopper+)**: distributed shared memory across a thread-block cluster.\n\nProducer/consumer pattern with these engines is what FA3 and FA4 codify.\n\n## 6. Ampere → Hopper → Blackwell evolution\n\n| Feature | Ampere (A100) | Hopper (H100/H200) | Blackwell (B200) |\n|---|---|---|---|\n| SMs | 108 | 132 | 148 |\n| Peak BF16 dense | 312 TFLOPs | 989 TFLOPs | ~2250 TFLOPs |\n| HBM | 40/80 GB HBM2e, 2.0 TB/s | 80 GB HBM3 3.35 TB/s / 141 GB HBM3e 4.8 TB/s | HBM3e ~8 TB/s |\n| Tensor core | 3rd gen (sync mma) | 4th gen (async **wgmma**) | 5th gen (**tcgen05**, 2-CTA MMA) |\n| Async copy | `cp.async` | **TMA** | TMA + TMEM |\n| Precision floor | INT8 / TF32 | FP8 | **FP4** (NVFP4/MXFP4), FP6 |\n| Special memory | — | Thread-block clusters + DSMEM | TMEM (256 KB/SM) |\n| Notable | first \"real\" tensor cores + BF16 | first async pipeline chip | first tensor memory chip |\n\nMicroarchitecture microbenchmark paper for Blackwell: arXiv 2512.02189 — B200 tensor cores measured 1.56× mixed-precision throughput and 42% better energy vs H200.\n\n## 7. Bandwidth vs compute — the numbers you must memorize\n\n**H100 SXM5:**\n- BF16 dense: 989 TFLOPs/s\n- HBM3: 3.35 TB/s\n- Ridge (BF16): 989e12 / 3.35e12 ≈ **295 FLOPs/byte** — need 295 FLOPs per byte moved to be compute-bound\n- FP8 dense: 1979 TFLOPs/s, ridge ≈ 590 FLOPs/byte\n\n**H200:** same GH100 die, memory upgrade → 141 GB HBM3e @ 4.8 TB/s. Ridge (BF16) drops to ~206 FLOPs/byte. **Memory-bound kernels get faster; compute-bound kernels don't.** That's why H200 wins on decode.\n\n**B200:** ~1.5× compute, ~2.4× BW vs H100. Ridge similar to H200 in BF16. Ridge for FP4/FP8 rises — low-precision kernels shift more toward memory-bound on Blackwell.\n\n**MI300X:** 192 GB HBM3 @ 5.3 TB/s, 750W. Bigger cache-coherent memory pool than any NVIDIA data-center chip — that's the reason Meta serves Llama 3.1 405B on it in production. CDNA3, matrix cores instead of tensor cores. Ecosystem: ROCm 7.\n\n**Consumer RTX 4090 (Ada, sm89):** 82.6 TFLOPs FP32, ~1.0 TB/s, no NVLink. Fine for learning + kernel dev; poor for real inference (24 GB, no FP8 tensor cores hardware).\n\n## 8. The roofline model — your first weapon\n\nDefinition:\n- **Arithmetic intensity** I = FLOPs / bytes-moved-from-HBM (for a given kernel + input shape)\n- **Peak perf** = min(Peak_FLOPs, I × Peak_BW)\n- **Ridge point** = Peak_FLOPs / Peak_BW (I where the two curves meet)\n\nExamples (H100 BF16, ridge ≈ 295):\n- **GEMM** with M=N=K=4096: intensity ≈ N ≈ 4096 → compute-bound. Your job: keep tensor cores fed.\n- **Elementwise / RMSNorm**: intensity ≈ O(1) → memory-bound. Your job: fuse it into the neighbor.\n- **Softmax** (naïve, 3 passes over the tensor): intensity ≈ 1–2 → memory-bound. Job: online softmax + fusion.\n- **Attention** (naïve, materialized N×N): reads/writes an N² tensor to HBM. FlashAttention avoids this by tiling and never materializing. Traffic goes from O(N²d) to O(N²d²/M) where M = SMEM size.\n\nRead: **Horace He, \"Making Deep Learning Go Brrrr From First Principles\"** — https://horace.io/brrr_intro.html — canonical rant that hammers this framework. Read it three times.\n\n## 9. Occupancy — the useful-but-overrated knob\n\nOccupancy = (active warps / SM) / (max warps / SM). Higher occupancy → more warps for the scheduler to hide latency with. But:\n\n- More occupancy = fewer registers/thread and less SMEM/block = smaller tiles.\n- Modern high-performance kernels (FA3/4, CUTLASS Hopper kernels) intentionally run at **low occupancy** with huge tiles + async pipelines to overlap. Occupancy is not the goal; **effective throughput** is.\n- Rule: use occupancy as a diagnostic (\"why is my kernel stalled?\"), not a target (\"maximize occupancy\").\n\n## 10. What to actually learn first\n\n1. **Say out loud, without looking:** thread / warp / warp scheduler / warp group / block / cluster / grid.\n2. **Draw the memory hierarchy** for H100 with bandwidths, from memory.\n3. **Compute the ridge point** for H100 BF16, H100 FP8, B200 BF16. Have the three numbers memorized.\n4. **Read Horace He's Brrr** post today. It's the mental model.\n5. Then move to `02_pmpp_path.md`.\n\n## References\n\n- Horace He, *Making Deep Learning Go Brrrr From First Principles*: https://horace.io/brrr_intro.html\n- NVIDIA H100 whitepaper: https://resources.nvidia.com/en-us-tensor-core\n- NVIDIA Blackwell architecture: https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/\n- PTX ISA (async operations): https://docs.nvidia.com/cuda/parallel-thread-execution/\n- Blackwell microbenchmark paper: arXiv 2512.02189\n- \"tcgen05 for dummies\" (Blackwell tensor core PTX walkthrough): https://gau-nernst.github.io/tcgen05\n