Warmup Guide — C++/CUDA & GPU Performance Engineering

Zero-to-senior primer for Phase 18. We start from "what is a GPU actually doing when it runs a kernel" — the warp in lockstep, the memory hierarchy, the SM scheduling warps to hide latency — and end with the four levers that move real kernel performance: coalescing, tiling, fusion, and the right precision. This phase deepens Phase 00's roofline from a model-level slide to a kernel-level tool, and adds the one calculation Phase 00 didn't: occupancy, and how to read the binding resource. By the end you can triage a slow kernel on a whiteboard.

Table of Contents

• Chapter 1: The GPU Execution Model — Thread, Warp, Block, Grid, SM
• Chapter 2: The Memory Hierarchy — Registers to HBM
• Chapter 3: Occupancy & Latency Hiding — Why GPUs Oversubscribe Warps
• Chapter 4: Memory Coalescing & Bank Conflicts — The #1 Lever
• Chapter 5: The Roofline at the Kernel Level (Deepening Phase 00)
• Chapter 6: Tiling & Blocking — Data Reuse and the Shared-Memory GEMM
• Chapter 7: Operator Fusion — FlashAttention, torch.compile, Triton
• Chapter 8: Tensor Cores & Mixed Precision
• Chapter 9: Streams, Asynchrony & CUDA Graphs — Launch Overhead
• Chapter 10: The Role of C++ — Custom Ops, Extensions, and Triple-Profiling
• Lab Walkthrough Guidance
• Success Criteria
• Interview Q&A
• References

Chapter 1: The GPU Execution Model — Thread, Warp, Block, Grid, SM

From zero. A CPU has a handful of fat cores tuned for latency: branch prediction, out-of-order execution, big caches — each core sprints through one thread. A GPU is the opposite: thousands of thin cores tuned for throughput. It does not try to make one thread fast; it tries to keep enormous numbers of threads in flight so there is always work to do while other threads wait on memory. To use it you express your problem as one function (a "kernel") run by thousands of threads at once, each thread handling a slice of the data.

The hierarchy. When you launch a kernel you launch a grid of blocks of threads:

  grid  ── many blocks (your whole problem)
   └─ block ── up to 1024 threads; shares SHARED MEMORY; can __syncthreads()
        └─ warp ── exactly 32 threads, executed in LOCKSTEP (SIMT)
             └─ thread ── one lane; has its own registers

Hardware side, the GPU has many Streaming Multiprocessors (SMs) — an A100 has 108. Each block is assigned to one SM and stays there; an SM runs many blocks (hence many warps) concurrently. The grid/block split is your logical decomposition; the warp and the SM are the physical units that decide performance.

The warp is the unit that matters. Threads are scheduled in groups of 32 called a warp, and all 32 lanes execute the same instruction at the same time on different data — this is SIMT (Single Instruction, Multiple Threads). The warp is the atom of execution: the scheduler issues one instruction per warp per cycle. Almost every GPU performance fact is really a fact about warps.

Warp divergence — the cost of branching. Because a warp is lockstep, what happens at if (threadIdx.x % 2 == 0) A(); else B();? The hardware cannot run A on 16 lanes while running B on the other 16. It runs A with the odd lanes masked off (idle), then runs B with the even lanes masked off. Both paths execute serially. A warp that splits k ways costs up to the sum of all k taken paths instead of one. This is warp_divergence_factor in the lab:

$$ \text{divergence factor} = #{\text{paths with a lane taking them}}. $$

A uniform warp (every lane takes the same path) has factor 1 — no penalty. A two-way split is up to 2×. This is why data-dependent branching inside a warp is a code smell, and why kernels are written to keep all 32 lanes on the same path (predication, sorting work so a warp is homogeneous, etc.).

Why this framing matters. When you later ask "why is this kernel slow," the answer is always about warps: warps stalled on memory (Chapter 2–4), too few warps to hide that stall (Chapter 3), or warps diverging (here). The execution model is the vocabulary for every diagnosis.

Common misconception. "A block is the parallel unit." No — the warp is. A block is a scheduling/sharing convenience (it gives you shared memory and __syncthreads()), but the instruction is issued per warp, and your block size matters mostly because it determines how cleanly it divides into warps and how many warps an SM can pack (Chapter 3).

Chapter 2: The Memory Hierarchy — Registers to HBM

From zero. A floating-point unit can do a multiply-add in ~a nanosecond, but reading a number from the GPU's main memory (HBM) takes hundreds of nanoseconds. If every operation waited on HBM, the math units would sit idle ~99% of the time. The entire art of GPU programming is keeping the data the math units need as close to them as possible. So you must know the levels:

Each step down is roughly an order of magnitude slower and bigger. Registers and shared memory are the only memories fast enough to keep up with the math units; everything else is the slow, shared bottleneck you are trying to touch as rarely as possible.

Registers are private per thread and the fastest storage — but they are a scarce, shared resource on the SM (Chapter 3): the more registers your kernel needs per thread, the fewer threads the SM can run at once. Shared memory is a small, programmer-managed scratchpad shared by all threads in a block; it is the staging area for the tiling trick (Chapter 6) and is physically the same SRAM as L1. L2 is a hardware cache for all SMs. HBM is the big off-chip memory where your tensors live — high bandwidth (TB/s) but high latency, and the resource the roofline is about.

        FAST/SMALL                                       SLOW/BIG
  registers → shared mem / L1 → L2 cache → HBM (global) → (PCIe → CPU RAM)
   ~1 cyc        ~20 cyc          ~200 cyc    ~500 cyc        ~µs
   private       per-block        per-GPU     per-GPU         host

Why this is the whole game. Every optimization in this phase is "move the working set up this ladder and keep it there": coalescing makes HBM reads efficient (Chapter 4), tiling stages data in shared memory to reuse it (Chapter 6), fusion keeps intermediates in registers instead of round-tripping HBM (Chapter 7). If you internalize the latency/bandwidth gaps, the optimizations stop being tricks and become obvious.

Common misconception. "More HBM bandwidth fixes a slow kernel." Sometimes — but usually the fix is to touch HBM less by reusing data on chip. A kernel that reads each input from HBM 100 times is not bandwidth-starved; it is wasting bandwidth, and tiling fixes it without faster memory.

Chapter 3: Occupancy & Latency Hiding — Why GPUs Oversubscribe Warps

The problem occupancy solves. From Chapter 2, an HBM read stalls a warp for hundreds of cycles. A CPU hides such stalls with caches and out-of-order execution. A GPU has a cheaper, brute-force trick: keep many warps resident on each SM, and when one warp stalls on memory, instantly switch to another that's ready. With enough warps in flight, the math units always have some warp to run, and the memory latency is hidden behind other warps' work. This is latency hiding by oversubscription, and it is the reason GPUs are built the way they are.

Occupancy is the metric: how many warps are actually resident on an SM, as a fraction of the maximum it can hold.

$$ \text{occupancy} = \frac{\text{active warps per SM}}{\text{max warps per SM}}. $$

Higher occupancy → more warps to hide latency behind → (usually) better tolerance of memory stalls.

What caps it — the four limits. An SM is a fixed bucket of resources, and a block consumes some of each. How many blocks (hence warps) fit is the minimum over four ceilings:

1. Register file. Each thread needs regs_per_thread registers; a block needs regs_per_thread × threads_per_block. With a 65,536-register file: blocks ≤ 65536 // (regs_per_thread × threads_per_block). A register-hungry kernel runs fewer warps — the classic occupancy killer.
2. Shared memory. A block reserves smem_per_block bytes; blocks ≤ smem_per_sm // smem_per_block. A kernel that grabs a big shared-memory tile is capped here.
3. Warp slots. The SM holds at most max_warps_per_sm warps: blocks ≤ max_warps_per_sm // warps_per_block.
4. Block slots. The SM holds at most max_blocks_per_sm blocks (e.g. 32) — tiny blocks hit this before anything else.

The binding (limiting) resource is whichever ceiling is lowest. Naming it is the whole point, because it tells you the fix:

This is exactly the lab's occupancy(...): it computes all four and returns active_warps, max_warps, the occupancy fraction, and the limiter string. A register-heavy kernel returns "registers"; a shared-mem-heavy one returns "shared_memory" — that is the occupancy soul test.

Worked example (the lab's A100-ish SM: 64 max warps, 65,536 regs, 64 KB smem, 32 block slots). A kernel of 256 threads/block (8 warps) using 128 registers/thread and no shared memory: register-limited blocks = 65536 // (128 × 256) = 2, so active_warps = 2 × 8 = 16 and occupancy = 16/64 = 0.25. Drop to 24 registers/thread and the warp slots bind instead (64 // 8 = 8 blocks), active_warps = 64, occupancy = 1.0.

The occupancy-vs-ILP tradeoff — the senior nuance. Higher occupancy is not always faster. A kernel can hide latency two ways: many warps (occupancy) or lots of independent work per thread (Instruction-Level Parallelism — issue several independent memory loads before using any of them). Volkov's famous result: a low-occupancy kernel with high ILP can beat a high-occupancy one, because the extra registers that lowered occupancy were spent on ILP that hid the latency anyway. So: occupancy is a means (hide latency), not the goal. ~50% is often plenty. Chase the limiter only until latency is hidden.

Common misconception. "Maximize occupancy." No — hide latency. Occupancy is one way; ILP is another. If you're already hiding latency at 50%, pushing to 100% by shrinking tiles or registers can make the kernel slower.

Chapter 4: Memory Coalescing & Bank Conflicts — The #1 Lever

From zero. The memory system does not serve one byte at a time; it serves fixed-size chunks — think a 128-byte segment (a cache line / a group of 32-byte sectors). When the 32 threads of a warp each issue a load, the hardware looks at all 32 addresses together and figures out how many distinct 128-byte segments they touch, then transfers exactly those segments. You pay for whole segments whether you use all of each one or not.

Coalesced access. If the 32 threads read 32 consecutive 4-byte words, that's 32 × 4 = 128 contiguous bytes = exactly one 128-byte transaction. Every byte transferred is used. Efficiency:

$$ \text{efficiency} = \frac{\text{useful bytes}}{\text{transferred bytes}} = \frac{32 \times 4}{1 \times 128} = 1.0. $$

This is coalescing, and it is the single highest-leverage thing in memory-bound GPU code: adjacent threads touch adjacent addresses → minimum transactions → full bandwidth.

Strided / scattered access. Now suppose thread i reads address i × stride × 4. With stride = 32, the 32 addresses are 128 bytes apart — each lands in its own 128-byte segment. You transfer 32 segments (32 × 128 = 4096 bytes) to use only 128 useful bytes:

$$ \text{efficiency} = \frac{32 \times 4}{32 \times 128} = \frac{1}{32} \approx 0.03. $$

You've thrown away ~97% of the bandwidth you paid for. This is coalescing_efficiency(stride) in the lab: stride 1 ≈ 1.0, stride 2 ≈ 0.5, stride 32 ≈ 1/32. That monotone fall-off is the coalescing soul test.

  stride 1 (coalesced):  |XXXXXXXX|  one transaction, all useful   → 1.00
  stride 2:              |X_X_X_X_|X_X_X_X_|  2 transactions, half  → 0.50
  stride 32:             |X_______|X_______|...  32 transactions   → 0.03

The canonical bug: AoS vs SoA. An array of structs Particle{x,y,z,...}[N] makes thread i read particles[i].x, which are sizeof(Particle) apart — strided, uncoalesced. A struct of arrays float x[N], y[N], z[N] makes thread i read x[i] — contiguous, coalesced. The same algorithm, an order of magnitude apart, decided purely by data layout. When a memory-bound kernel is "mysteriously slow," check coalescing first; it is usually an AoS that wants to be SoA.

Bank conflicts (the shared-memory analogue). Shared memory is split into 32 banks; a warp can read 32 words in one cycle only if they hit 32 distinct banks. If two lanes hit the same bank (e.g. a column access in a [32][32] tile, where every element of a column is in the same bank), the accesses serialize — an N-way bank conflict is N× slower. The standard fix is padding (__shared__ float tile[32][33];) so columns skew across banks. Same idea as coalescing, one level up the hierarchy.

Why this is the #1 lever. Decode, attention, and most LLM kernels are memory-bound (Chapter 5), so their speed is their effective bandwidth, and coalescing is effective bandwidth. A coalescing fix is often the largest, cheapest win available — no algorithm change, just a layout or index fix.

Common misconception. "Coalescing is automatic / the compiler handles it." The compiler emits the loads; you decide the addresses through your indexing and data layout. The compiler cannot turn an AoS into an SoA.

Chapter 5: The Roofline at the Kernel Level (Deepening Phase 00)

Recall Phase 00. A GPU has two peak rates — peak compute (FLOP/s) and memory bandwidth (bytes/s). For any workload, arithmetic intensity I = FLOPs / bytes_moved (FLOP/byte) decides which one binds, via the roofline:

$$ \text{attainable FLOP/s} = \min(\text{peak FLOP/s},; \text{bandwidth} \times I), \qquad I_{\text{ridge}} = \frac{\text{peak}}{\text{bandwidth}}. $$

Phase 00 applied this to a whole model (decode I≈1 → memory-bound). Phase 18 applies the same formula to a single kernel, which is where it becomes an optimization tool rather than a slide.

 FLOP/s
   peak ┤        ________________   ← compute-bound (flat roof)
        │       /        ● big GEMM (tiled): near the ridge, compute-bound
        │      /  slope = bandwidth
        │     /
        │  ● decode / elementwise: I≈1, far left → MEMORY-BOUND
        └───┴──────────────────────► arithmetic intensity (FLOP/byte)
          ridge = peak/bw ≈ 156 on an A100

The two regimes, per kernel.

• Left of the ridge (I < ridge) → memory-bound: you finish the math before the data arrives. The lever is bytes: coalesce (Chapter 4), tile (Chapter 6), fuse (Chapter 7), use a smaller dtype. More FLOP/s does nothing. Element-wise ops, decode, softmax, layernorm, and unfused attention live here.
• Right of the ridge (I ≥ ridge) → compute-bound: the math units are saturated. The lever is better math: tensor cores, mixed precision (Chapter 8), fewer FLOPs. Faster memory does nothing. Large dense GEMMs and prefill live here.

This is the lab's roofline_attainable, ridge_point, and is_compute_bound — the same min() as Phase 00, now used to classify this kernel and pick its lever.

Arithmetic intensity of a GEMM, and why tiling raises it. A naive C[M,N]=A[M,K]·B[K,N] does 2·M·N·K FLOPs. Its intensity depends entirely on how many times it re-reads A and B. If you could read each element exactly once, bytes = (M·K + K·N + M·N)·b and I is high (gemm_arithmetic_ intensity). But a naive kernel re-reads a row of A for every column of N — it moves far more bytes, so its effective intensity is low and it's memory-bound. Tiling moves the kernel rightward on the roofline by cutting the bytes (Chapter 6). The lab models both the idealized intensity (gemm_arithmetic_intensity) and the actual traffic of a tiled schedule (tiled_gemm_hbm_traffic), and the gap between them is the optimization opportunity.

The punchline for LLMs. Decode reads every weight (~2 bytes) to do ~2 FLOPs → I≈1, far left of 156 → hopelessly memory-bound (the P00 result). So the inference levers are quantize (fewer bytes), batch (reuse the weight read across sequences → higher I), and fuse — never "buy more FLOPs."

Common misconception. "nvidia-smi says 100%, so we're saturated." Utilization = "a kernel was running," not "the math units were busy." A memory-bound kernel shows 100% util while doing single- digit % of peak FLOPs (low MFU). The roofline chart in Nsight, not the util meter, tells the truth — exactly the Phase 00 lesson, now provable per kernel.

Chapter 6: Tiling & Blocking — Data Reuse and the Shared-Memory GEMM

The problem. A naive GEMM is the canonical memory-bound disaster: to compute one output element C[i][j] you read a whole row of A and a whole column of B from HBM, and you do this for every one of M·N outputs — each A and B element is dragged from HBM O(N) and O(M) times. The math is fast; the kernel spends all its time waiting on HBM.

The fix: tiling (a.k.a. blocking). Chop the output into tile×tile blocks. To compute one output tile, load a tile×tile block of A and of B into shared memory once, then let every one of the tile×tile threads reuse those loaded values tile times before moving to the next K-block. Each HBM byte now feeds tile multiply-adds instead of one. Total HBM traffic:

$$ \text{reads} = \frac{2,M,N,K}{\text{tile}} \cdot b, \qquad \text{writes} = M,N \cdot b ;;\Rightarrow;; \text{traffic} \propto \frac{1}{\text{tile}}. $$

That 1/tile is the lab's tiled_gemm_hbm_traffic: a bigger tile reuses more, so it moves less data. On a 4096³ fp16 GEMM the naive schedule (tile=1) moves ~275 GB; tile=32 moves ~8.6 GB (~32× less); tile=128 moves ~2.2 GB (~126× less). That is the tiling soul test — and it is the single most important kernel optimization, because it converts a memory-bound matmul into a compute-bound one (it slides the kernel right across the ridge in Chapter 5).

  naive: read A row & B col from HBM for EVERY output element   → traffic ∝ M·N·K
  tiled: load TILE×TILE blocks to shared mem, reuse TILE times  → traffic ∝ M·N·K / tile
                       ┌────────┐   reused by every thread
        HBM ──load──▶  │ shared │   in the block, TILE times
                       │  tile  │   ──▶ registers ──▶ FMA
                       └────────┘

What limits the tile (the tradeoff). Bigger tile = less traffic, but the tile lives in shared memory and its partial sums in registers — both are the scarce SM resources from Chapter 3. Push the tile too far and you blow the shared-memory or register budget and occupancy collapses, so you can't hide the latency anymore. The optimal tile balances reuse (Chapter 6) against occupancy (Chapter 3) — the central tension of GEMM tuning, and exactly why CUTLASS/cuBLAS autotune tile shapes per GPU.

Common misconception. "Tiling is a matmul-only trick." It's the general principle of blocking for cache/scratchpad reuse — it shows up in convolutions, attention (FlashAttention tiles the score matrix), stencils, and any kernel that revisits data.

Chapter 7: Operator Fusion — FlashAttention, torch.compile, Triton

The problem. A chain of element-wise ops — say y = gelu(x); z = y * w; out = z + b — is, in a naive framework, three kernels. Each one reads its input from HBM and writes its output back. So a 3-op chain does 2 × 3 = 6 HBM passes over the data, even though the math per element is trivial. Element-wise ops have I ≈ tiny, so they are deeply memory-bound (Chapter 5): the time is entirely the HBM round-trips, and most of those round-trips are writing an intermediate only to immediately read it back.

The fix: fusion. Run the whole chain in one kernel: read x from HBM once, compute gelu → multiply → add entirely in registers, write out once. Now it's 2 HBM passes regardless of how many ops are in the chain:

$$ \frac{\text{unfused}}{\text{fused}} = \frac{2 \cdot n_{\text{ops}} \cdot \text{elems} \cdot b} {2 \cdot \text{elems} \cdot b} = n_{\text{ops}}. $$

That ratio is the lab's fused_vs_unfused_traffic — fusing a 5-op chain cuts HBM traffic ~5×, and for a bandwidth-bound chain that's a ~5× speedup. Fusion and tiling are the same idea (pay HBM less often); tiling reuses input data on-chip, fusion keeps intermediate data on-chip.

FlashAttention — the canonical fused, IO-aware kernel. Standard attention computes S = QKᵀ (an N×N score matrix), softmaxes it, then O = softmax(S)·V. The naive version writes the entire N×N matrix to HBM and reads it back — for long sequences that matrix is huge and the kernel is memory-bound on materializing it. FlashAttention tiles Q, K, V into blocks, computes attention block-by-block in shared memory using the online softmax trick (running max + running normalizer so it never needs the whole row at once), and never writes S to HBM at all. It is tiling (Chapter 6) + fusion (here) + numerical care, and it's why long-context attention became practical. When an interviewer asks "what does FlashAttention actually do," the answer is: it's an IO-aware fused kernel that keeps the score matrix on chip.

What torch.compile and Triton do. You rarely hand-fuse anymore:

• torch.compile traces your model, finds fusible op chains, and generates fused kernels (via TorchInductor, which emits Triton) — automatic fusion for the common cases.
• Triton is a Python DSL for writing GPU kernels at the block level: you write a @triton.jit function over tiles, and Triton handles the thread-level details, coalescing, and autotuning. It's the middle path — most of CUDA's speed for a fraction of the C++ effort, and what most custom LLM kernels (including many FlashAttention variants) are now written in.

Common misconception. "Fusion is a micro-optimization." For memory-bound element-wise/normalize chains it's often the whole win, because the kernel was never compute-limited — it was paying HBM for intermediates it never needed to store.

Chapter 8: Tensor Cores & Mixed Precision

What a tensor core is. A regular GPU core does one FMA per cycle. A tensor core is a dedicated unit that does a small matrix-multiply-accumulate (e.g. 16×16 tiles) per instruction — an order of magnitude more FLOP/s than the regular cores, for matmuls in the right precision. An A100's ~312 TFLOP/s bf16 number is a tensor-core number; the non-tensor-core fp32 rate is ~5–6× lower. If your GEMM isn't using tensor cores, you're leaving most of the chip on the table.

The catch: precision. Tensor cores want lower-precision inputs — fp16, bf16, tf32, fp8 — and accumulate in fp32. So the speed is unlocked by mixed precision: store/multiply in 16-bit (or 8), accumulate in 32-bit to keep numerical sanity. bf16 (same exponent range as fp32, fewer mantissa bits) is the modern default for training because it rarely overflows; fp16 has more precision but a narrow range (needs loss scaling); fp8 (H100+) pushes further for inference and some training.

Why it ties to the roofline. Mixed precision moves you up and right on the roofline at once: up because tensor cores raise the peak-FLOP/s ceiling, right because 16-/8-bit data halves/quarters the bytes moved (higher I). It's the rare optimization that helps both regimes — which is why it's the first thing every serious GEMM/attention kernel does, and why P06 (quantization) and this phase are two sides of one coin.

The constraint. Tensor cores need matrix dimensions that are multiples of 8/16 (and like larger tiles) to hit peak — odd shapes silently fall back to slow paths. "Pad to a multiple of 8" is real, free performance advice.

Common misconception. "Mixed precision just saves memory." It saves memory and unlocks the tensor cores — the speedup is usually the bigger deal, and it's why bf16 training is ~2–3× faster, not just smaller.

Chapter 9: Streams, Asynchrony & CUDA Graphs — Launch Overhead

The hidden tax: kernel-launch overhead. Every kernel launch costs the CPU a few microseconds to set up and dispatch. For a big kernel that runs for milliseconds, irrelevant. But decode emits one token at a time, and each token is dozens to hundreds of tiny kernels (one per layer per op), each running for microseconds. Now the few-microsecond launch overhead is a large fraction of the work — the GPU sits idle between kernels waiting for the CPU to launch the next one. Launch overhead bites decode hardest, precisely because decode's kernels are tiny and numerous.

Streams & asynchrony. A CUDA stream is an ordered queue of operations; the CPU enqueues kernels/copies asynchronously and moves on. Independent work on different streams can overlap — classically, copy data for batch n+1 (a memory-transfer stream) while computing batch n (a compute stream), so the PCIe copy hides behind compute. Asynchrony is how you keep the GPU fed.

CUDA Graphs — the launch-overhead killer. Instead of the CPU launching the same 200 kernels every decode step, you capture that sequence once into a graph and then replay the whole graph with a single launch. The CPU overhead collapses from "200 launches" to "1 replay," and the GPU stops starving between kernels. This is a major decode-latency optimization (vLLM, TensorRT-LLM, and torch.compile's mode="reduce-overhead" all use CUDA Graphs) — and it only matters because of the launch-overhead problem above.

  without graphs:  CPU: launch k1 ─▶ launch k2 ─▶ ... (GPU idle in the gaps)
  with graphs:     CPU: replay(graph) ───▶  GPU runs k1..kN back-to-back, no gaps

Common misconception. "Launch overhead is negligible." For prefill/training (big kernels), yes; for decode (tiny kernels), it can be a third of your latency, which is why CUDA Graphs are standard in production decode and why "why is my single-stream decode slow even though the GPU is fast" often answers to launch overhead, not compute.

Chapter 10: The Role of C++ — Custom Ops, Extensions, and Triple-Profiling

Where C++ enters. PyTorch is Python on top of C++ (ATen) on top of CUDA kernels. Most of the time you live in Python and the framework's kernels are fine. You drop to C++/CUDA when:

• an op doesn't exist or fuses poorly and is a measured hot spot,
• you need a custom kernel (a novel attention variant, a fused norm, a quantized matmul),
• you need control the framework won't give you (specific tiling, tensor-core usage, smem layout).

The PyTorch C++/CUDA extension is the path: write the kernel in .cu, a thin .cpp binding (via torch::Tensor), and build it with torch.utils.cpp_extension (load(...) for JIT, or a setup.py with CUDAExtension). Now import myext; myext.my_op(t) calls your kernel as a normal torch op, autograd-compatible if you register a backward. That's how every custom-op library (FlashAttention, xFormers, bitsandbytes) ships.

The three rungs — pick the cheapest that hits your target:

The senior judgment is not "always write CUDA" — it's "reach for the lowest rung that clears the bar," and being able to climb to the top rung when needed.

Profiling — measure, never guess. Two tools:

• Nsight Systems (nsys) — the timeline: kernel launches, gaps, stream overlap, CPU↔GPU. Use it to find launch overhead, serialization, and idle gaps (Chapter 9).
• Nsight Compute (ncu) — per-kernel deep dive: the roofline chart, occupancy and its limiter, memory coalescing (sectors/request), warp-state stalls, tensor-core usage. This is where every number in this lab shows up for real. ncu --set roofline ./k puts your kernel on the chart.
• nvidia-smi is not a profiler. Its "utilization" means "a kernel ran," not "the math units were busy" — the Phase 00 warning, repeated because people keep believing it. Use MFU (achieved/peak FLOPs) and the roofline for the truth.

Common misconception. "Optimize, then measure." Reverse it. Profile first to find the one binding bottleneck (a roofline placement, an occupancy limiter, an uncoalesced access, a launch- overhead gap), fix that, re-profile. Guessing optimizations on a GPU is how you spend a week speeding up something that wasn't the bottleneck.

Lab Walkthrough Guidance

The lab (lab-01-gpu-roofline-occupancy) turns these chapters into a runnable model — no GPU required. Suggested order (matches the file):

1. Roofline (roofline_attainable, ridge_point, is_compute_bound) — Chapter 5. Same min() as Phase 00; is_compute_bound is intensity >= ridge (the ridge itself counts as compute-bound). Tie the test numbers to the A100 ridge of 156.
2. GEMM intensity (gemm_arithmetic_intensity) — Chapter 5/6. FLOPs 2·M·N·K over the naive bytes; note how tiling (next) lowers the bytes and raises effective intensity.
3. Occupancy (occupancy) — Chapter 3. The heart of the lab: compute warps/block, then the four limits, take the min, and return the limiter string. The register-limited vs shared-mem-limited tests are the soul test — get the limiter right.
4. Coalescing (coalescing_efficiency) — Chapter 4. Build the set of touched transaction_bytes-aligned segments across the 32 lanes; useful / transferred. Stride 1 → 1.0; stride 32 → 1/32.
5. Tiling (tiled_gemm_hbm_traffic) — Chapter 6. ceil-divide for the tile counts; A+B reads ∝ 1/tile, C writes once. A larger tile → fewer bytes (the tiling soul test).
6. Fusion (fused_vs_unfused_traffic) — Chapter 7. 2·n_ops passes vs 2; ratio = n_ops.
7. Divergence (warp_divergence_factor) — Chapter 1. Validate fractions sum to 1.0; factor = count of nonzero paths.

Run red → green: pytest test_lab.py -v, then LAB_MODULE=solution pytest test_lab.py -v, then python solution.py for the worked numbers. Then build the real CUDA kernels in the lab README's Extensions on a GPU.

Success Criteria

• All tests pass against your lab.py and against LAB_MODULE=solution.
• You can draw thread → warp (32, SIMT) → block → grid → SM and explain warp divergence from it.
• You can recite the memory hierarchy (registers → shared/L1 → L2 → HBM) with the latency/size ordering and say which optimization targets which level.
• You can compute an occupancy by hand for a (threads, regs, smem) kernel and name the binding resource, and say what you'd change to raise it (and why you might not).
• You can explain why coalescing_efficiency(1) ≈ 1.0 and coalescing_efficiency(32) ≈ 1/32, and why AoS→SoA is the classic coalescing fix.
• You can explain why tiling makes HBM traffic ∝ 1/tile, and what stops you from using an arbitrarily large tile (occupancy).
• You can explain fusion's n_ops× ratio and describe FlashAttention as tiling + fusion that keeps the score matrix off HBM.
• You can take any "this kernel is slow" prompt and produce: roofline placement → occupancy limiter → coalescing estimate → the one highest-leverage fix.

Interview Q&A

• "Thread vs warp vs block vs SM — what actually executes?" — The warp (32 threads, SIMT lockstep) is the unit of execution; a block is a scheduling/sharing group (shared memory, __syncthreads) pinned to one SM; the grid is your whole problem; the SM is the physical core that runs many warps to hide latency.
• "Is this kernel compute- or memory-bound? Prove it." — Compute I = FLOPs/bytes, compare to the ridge peak/bw (~156 on an A100). Decode I≈1 → memory-bound (batch/quantize/fuse); a big GEMM I≈500 → compute-bound (tensor cores). Cite the roofline, not nvidia-smi.
• "Your occupancy is 25%. What's wrong and how do you fix it?" — Find the limiter. If register-limited, cut register pressure (__launch_bounds__, recompute vs cache) or accept it for ILP; if smem-limited, shrink the tile; if you're warp-slot-limited you're already saturated. And remember occupancy is a means to hide latency, not a goal.
• "What's the single biggest memory optimization?" — Coalescing. Make adjacent threads touch adjacent addresses so a warp's loads pack into the fewest transactions; the classic fix is AoS→SoA. Strided access wastes bandwidth proportional to the stride.
• "Why does tiling speed up GEMM, and what limits the tile size?" — It loads a block into shared memory and reuses each element tile times, cutting HBM traffic ∝ 1/tile and pushing the kernel across the ridge to compute-bound. The tile is capped by shared-memory/register budget, beyond which occupancy collapses — the reuse-vs-occupancy tradeoff.
• "What does FlashAttention actually do?" — It's an IO-aware fused, tiled attention kernel: it tiles Q/K/V, uses an online (running-max) softmax, and never materializes the N×N score matrix in HBM, turning a memory-bound op into an on-chip one. Fusion + tiling + numerical care.
• "What are tensor cores and how do I use them?" — Dedicated matrix-multiply-accumulate units that deliver the headline FLOP/s, but only for low-precision inputs (fp16/bf16/tf32/fp8, fp32 accumulate) and dims that are multiples of 8/16. They move you up and right on the roofline.
• "Why is single-stream decode slow even on a fast GPU?" — Often kernel-launch overhead: decode is many tiny kernels per token, so the few-µs launch cost dominates and the GPU starves between kernels. CUDA Graphs (replay the captured sequence in one launch) fix it; so does bigger batching.
• "When do you write CUDA C++ vs Triton vs torch.compile?" — Lowest rung that clears the bar: torch.compile first (free fusion), Triton for a custom/fused kernel without C++, CUDA C++ extension only when you need control or the absolute peak. Senior judgment is restraint, not always reaching for CUDA.
• "nvidia-smi shows 100% — are we optimal?" — No. Utilization = "a kernel ran," not "FLOPs busy." Use MFU and the Nsight roofline; a memory-bound kernel is 100% util at low MFU.
• "What is warp divergence and how do you avoid it?" — Lanes taking different branches force the warp to run all paths serially (masked). Avoid by making branches uniform across a warp (sort work so a warp is homogeneous, use predication, restructure data-dependent control flow).
• "How do you find the bottleneck on a real kernel?" — Profile, don't guess: nsys for the timeline (launch overhead, gaps, overlap), ncu for per-kernel roofline/occupancy/coalescing/ warp-stalls. Fix the one binding cause, re-profile.

References

• NVIDIA, CUDA C++ Programming Guide — the execution model, memory hierarchy, coalescing, shared memory/bank conflicts, streams, and CUDA Graphs (the primary source).
• Williams, Waterman, Patterson, Roofline: An Insightful Visual Performance Model (2009) — the roofline this phase deepens from Phase 00.
• Dao et al., FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness (2022) and FlashAttention-2 (2023) — the canonical fused, tiled, IO-aware kernel.
• Tillet et al., Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations (2019) — the block-level kernel DSL; and the OpenAI Triton docs.
• Simon Boehm, How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance — the best step-by-step on coalescing, tiling, and occupancy in practice.
• V. Volkov, Better Performance at Lower Occupancy (GTC 2010) — the occupancy-vs-ILP tradeoff.
• Kirk & Hwu, Programming Massively Parallel Processors — the standard textbook; and the GPU MODE lecture series / Discord for modern kernel engineering.
• NVIDIA Nsight Compute & Nsight Systems docs — the roofline, occupancy, and memory-analysis sections that every number in this lab maps to on real hardware.
• PyTorch docs — Custom C++ and CUDA Extensions, and torch.compile / TorchInductor.