Hitchhiker's Guide — C++/CUDA & GPU Performance Engineering

The compressed practitioner tour. If WARMUP.md is the professor, this is the senior who leans over and says "here's what you actually need to remember when a kernel is slow."

The 30-second mental model

A GPU runs your kernel as warps (32 threads in lockstep, SIMT) packed onto SMs, and hides the brutal HBM latency by keeping many warps in flight (occupancy). Speed is decided by the memory hierarchy — registers → shared/L1 → L2 → HBM — and the game is touch HBM as rarely as possible. Put any kernel on the roofline (min(peak, bw·I), ridge = peak/bw): left of the ridge it's memory-bound (fix bytes: coalesce, tile, fuse, lower precision), right of it it's compute-bound (fix math: tensor cores, mixed precision). Decode is I≈1 → memory-bound, always. nvidia-smi 100% ≠ efficient — read the roofline.

The numbers to tattoo on your arm

Back-of-envelope one-liners

# "Compute- or memory-bound?"
I = FLOPs / bytes;  ridge = peak/bw (~156 A100);  I < ridge → memory-bound, else compute-bound

# "Occupancy of 256-thread, 128-reg/thread kernel on a 65536-reg SM?"
blocks = 65536 // (128*256) = 2  → 2*8 = 16 warps / 64 max = 25% (REGISTER-limited)

# "Coalescing efficiency at stride 8 (4B words, 128B tx)?"
~ 32 lanes scatter across ~8 transactions → 128 useful / 1024 transferred ≈ 0.125

# "How much does tiling save a 4096³ GEMM?"
tile 1 → ~275 GB HBM;  tile 32 → ~8.6 GB (~32x);  tile 128 → ~2.2 GB (~126x)

# "Fuse a 5-op element-wise chain?"
2*5 passes → 2 passes = 5x less HBM traffic (≈ 5x faster, it was bandwidth-bound)

The framework one-liners (where these live in real tools)

# Try fusion for free first:
model = torch.compile(model)                    # TorchInductor → Triton, auto-fuses
model = torch.compile(model, mode="reduce-overhead")   # also captures CUDA Graphs

# A custom fused kernel without writing CUDA C++:
import triton, triton.language as tl
@triton.jit
def my_kernel(x_ptr, ...): ...                  # block-level, autotuned tiles

# Drop to CUDA C++ when you must:
from torch.utils.cpp_extension import load
ext = load(name="myop", sources=["op.cpp", "kernel.cu"])  # JIT-build a torch op

# See registers/smem the compiler chose (the occupancy inputs):
#   nvcc --ptxas-options=-v   →  "Used N registers, M bytes smem"
# Profile for real:
#   ncu --set roofline ./k    →  roofline chart, occupancy + LIMITER, sectors/request
#   nsys profile ./app        →  timeline: launch overhead, gaps, stream overlap

War stories

• The 10× coalescing win. A particle kernel ran at ~8% of HBM bandwidth. The struct was Particle{x,y,z,vx,...} — an array of structs, so each thread's .x reads were strided and uncoalesced. Flip to struct-of-arrays (float x[N], y[N], ...) and it coalesced to ~full bandwidth. Same math, one data-layout change, an order of magnitude. Check coalescing first.
• The "max occupancy" regression. Someone shrank tiles to hit 100% occupancy and the GEMM got slower — the smaller tiles killed data reuse and ILP, so it became HBM-bound again. Occupancy is a means (hide latency), not the goal. ~50% with big reuse beat 100% with none.
• The decode that wouldn't speed up. Single-stream decode was slow on a fast GPU; the profiler showed the GPU idle in tiny gaps between hundreds of microsecond kernels — launch overhead. CUDA Graphs (capture once, replay) and bigger batching, not a faster GPU, fixed it.
• The nvidia-smi trap, again. "GPU's at 100%, we're maxed out." It was a memory-bound kernel at ~5% MFU — 100% util means a kernel ran, not that FLOPs were busy. The roofline showed it sitting far left of the ridge; the fix was fusion, not a bigger card.
• The bank-conflict mystery. A tiled transpose was 32× slower than expected — every thread hit the same shared-memory bank reading a column. Pad the tile ([32][33]) so columns skew across banks; conflict gone.

Vocabulary (rapid-fire)

• Warp — 32 threads executed in lockstep (SIMT); the real unit of execution.
• SM — streaming multiprocessor; runs many warps to hide latency.
• Occupancy — active warps ÷ max warps per SM; capped by regs / smem / warp slots / block slots.
• Limiter — the lowest of those four ceilings; the lever you pull.
• Coalescing — a warp's loads packing into the fewest memory transactions (contiguous = best).
• Bank conflict — shared-memory accesses hitting the same bank; serializes a warp.
• Arithmetic intensity — FLOPs per byte moved; where you sit on the roofline.
• Tiling/blocking — stage a block in shared memory and reuse it; traffic ∝ 1/tile.
• Fusion — one kernel for an op chain; HBM only at the boundaries (FlashAttention).
• Tensor core — matrix-multiply-accumulate unit; the headline FLOP/s, low precision only.
• MFU — achieved ÷ peak FLOPs; the honest efficiency number (not nvidia-smi util).
• CUDA Graph — capture a kernel sequence, replay in one launch; kills launch overhead.
• Triton — Python DSL for block-level kernels; the middle path between PyTorch and CUDA C++.

Beginner mistakes

• Thinking the block is the parallel unit — it's the warp (32, lockstep).
• Ignoring coalescing / shipping an array-of-structs that wants to be a struct-of-arrays.
• Chasing 100% occupancy instead of hiding latency (occupancy-vs-ILP).
• Reporting occupancy without the limiter — useless without "which resource."
• Forgetting tensor-core dim multiples (8/16) and silently falling back to slow paths.
• Treating nvidia-smi utilization as efficiency (use MFU + roofline).
• Optimizing before profiling — speeding up something that wasn't the bottleneck.
• Reaching for CUDA C++ when torch.compile/Triton would've done it in an afternoon.
• Forgetting launch overhead on decode — many tiny kernels, CPU-bound between them.

The one thing to take away

Before you touch a slow kernel, place it on the roofline, compute its occupancy and name the limiter, estimate its coalescing, and pick the one highest-leverage fix (coalesce / tile / fuse / cut registers / graph). Then profile to confirm. If you can do that on a whiteboard with no GPU in the room, you own the bottom of the stack — and that's the role this whole curriculum was building toward.