Phase 2 — GPU Architecture & Kernel Programming¶
Duration target: ~10 weeks (Months 3–5, running in parallel with Phase 3 in the second half). Prerequisite: Phase 1 (transformer math, from-scratch attention on CPU/GPU tensors in PyTorch).
Why this phase exists¶
An inference engineer who cannot read a kernel, name the memory hierarchy, or pick a tile size is a PyTorch application developer, not an inference engineer. Every performance conversation — batching, quantization, attention variants, fusion — bottoms out in the same three questions:
Am I memory-bound or compute-bound? (arithmetic intensity vs roofline)
Where is data living right now? (HBM → L2 → SMEM → RF)
Are the tensor cores busy? (occupancy, warp scheduling, async pipelines)
If those three questions do not have crisp answers for a given kernel, you are guessing. Phase 2 makes them reflexive.
Mental model¶
HBM (slow, huge) <— 3.35 TB/s (H100), ~8 TB/s (B200)
│
L2 (fast, shared, ~50 MB)
│
┌────┴────────────────────────────────────┐
│ SM 1 SM 2 ... SM N (132 H100 / 148 B200) │
│ ┌────────────────────────────────────┐ │
│ │ SMEM (~228 KB/SM H100) │ │
│ │ TMEM (256 KB/SM Blackwell only) │ │
│ │ Registers (65,536 x 32-bit / SM) │ │
│ │ 4 warp schedulers, tensor cores │ │
│ └────────────────────────────────────┘ │
└──────────────────────────────────────────────────┘
Every optimization is a variation of: move data less, keep tensor cores fed.
The three lenses you must acquire¶
Roofline lens — for any kernel, compute arithmetic intensity (FLOPs / bytes moved) and place it on the roofline. Below ridge = memory-bound (fuse, tile, reuse). Above ridge = compute-bound (better MMA, occupancy, async).
Tile lens — every fast kernel is a nested tiling: block-tile in SMEM, warp-tile in registers/tensor cores, thread-tile in registers. You will name every tile.
Pipeline lens — modern kernels are producer/consumer async pipelines (TMA loads → wgmma → epilogue). “Optimize” almost always means “overlap something with something else.”
Exit criteria (do not leave this phase until all pass)¶
SGEMM ladder complete: Your CUDA SGEMM hits ≥ 70–80% of cuBLAS on fp32 on your target GPU (RTX 4090 / L4 / H100 as available), matching Simon Boehm’s kernel 8–10 range.
Fused softmax in Triton from memory — you can write, autotune, and benchmark it without looking at the tutorial.
Roofline placement: given a kernel + GPU spec, you can compute the ridge point and state whether the kernel is memory- or compute-bound within 30 seconds.
Nsight Compute fluency: for any kernel you write, you can pull achieved memory throughput, SM efficiency, tensor-core utilization, and identify the top bottleneck.
You know these numbers cold: H100 SXM = 989 TFLOPs BF16 dense / 3.35 TB/s / 132 SMs / ridge ≈ 295 FLOPs/byte BF16. B200 ≈ 1.5× compute, ≈ 2.4× BW. MI300X = 192 GB / 5.3 TB/s.
You have opinions: you can articulate when to reach for Triton vs CUTLASS vs raw CUDA, and why FA4 is written in CuTe-DSL not Triton.
File map for this phase¶
File |
Purpose |
|---|---|
|
The physical machine: SMs, warps, tensor cores, memory hierarchy, Ampere→Hopper→Blackwell, roofline, numbers per generation |
|
PMPP 4th ed chapter map with inference-specific exercise picks |
|
The GPU MODE community + KernelBench/KernelBot leaderboards + lecture prio |
|
Annotated Simon Boehm ladder — each step, expected speedup, Nsight metric to verify |
|
Triton tutorial order + modern (2025–2026) considerations + Gluon note |
|
CUTLASS/CuTe learning path + when to invest + CuTe-DSL note |
|
Nsight Compute + Systems + torch.profiler workflow, benchmark hygiene |
|
ROCm/HIP + Apple MLX primer (awareness-level) |
|
Concrete Phase 2 kernels to write with acceptance criteria |
The “no you don’t need to know” list¶
To keep you from rabbit-holing: for inference engineering you do not need mastery of graphics pipelines, CUDA driver API (runtime API is fine), texture memory, cooperative groups edge cases, PTX assembly (read it, don’t write it), OptiX, or MIG partitioning. Skip.
What comes next¶
Phase 3 (../04_attention/) rides directly on this foundation — you cannot understand FlashAttention without the memory hierarchy, and you cannot write it in Triton without the tile lens.