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:

  1. Am I memory-bound or compute-bound? (arithmetic intensity vs roofline)

  2. Where is data living right now? (HBM → L2 → SMEM → RF)

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

  1. 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).

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

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

01_gpu_hardware_model.md

The physical machine: SMs, warps, tensor cores, memory hierarchy, Ampere→Hopper→Blackwell, roofline, numbers per generation

02_pmpp_path.md

PMPP 4th ed chapter map with inference-specific exercise picks

03_gpu_mode.md

The GPU MODE community + KernelBench/KernelBot leaderboards + lecture prio

04_sgemm_ladder.md

Annotated Simon Boehm ladder — each step, expected speedup, Nsight metric to verify

05_triton_path.md

Triton tutorial order + modern (2025–2026) considerations + Gluon note

06_cutlass_cute.md

CUTLASS/CuTe learning path + when to invest + CuTe-DSL note

07_profiling_mastery.md

Nsight Compute + Systems + torch.profiler workflow, benchmark hygiene

08_amd_apple_awareness.md

ROCm/HIP + Apple MLX primer (awareness-level)

09_projects.md

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.