🛸 Hitchhiker's Guide — Phase 02: CUDA Programming

Read this if: you understand the Phase 01 machine model but haven't shipped CUDA. This is the field guide: idioms, numbers, foot-guns, and the review standards a GPU platform team should hold.

0. The 30-second mental model

CUDA = C++ where one function runs N times concurrently, N = grid × block. You manage three things: indices (which element am I?), memory (is my warp's access coalesced? is reuse staged in shared memory?), and asynchrony (launches return immediately; streams order work; events time and connect it). Performance work = move the kernel toward the correct roof from Phase 01's roofline, then stop.

1. The idioms you'll type forever

// ceil-div launch + tail guard — THE pattern
int t = 256, b = (n + t - 1) / t;
k<<<b, t>>>(d_x, n);                 // d_x from cudaMalloc, never a host ptr
// error discipline
CUDA_CHECK(cudaGetLastError());      // launch errors (sync)
CUDA_CHECK(cudaDeviceSynchronize()); // kernel errors (async) — debug builds
// honest timing
cudaEventRecord(t0); k<<<...>>>(...); cudaEventRecord(t1);
cudaEventSynchronize(t1); cudaEventElapsedTime(&ms, t0, t1);

Block size: multiple of 32, default 128/256, decided by measurement not vibes. 3D grids are indexing sugar — the hardware sees a flat list of blocks.

2. Memory decisions, ranked by frequency of regret

1. Sync copies in hot loops (PCIe is 1/50th of HBM): batch, pin, overlap.
2. Uncoalesced layouts: AoS structs, column walks of row-major data, threadIdx.y-fastest indexing. Fix = SoA + adjacent-thread-adjacent-address. ncu smoking gun: sectors/request ≈ 32 (want ~4).
3. Unified memory in serving paths: page-fault storms. Managed memory is a prototyping tool; production transfers are explicit.
4. Raw cudaMalloc/Free per request: device-synchronizing and fragmenting — use cudaMallocAsync pools or a caching allocator (you'll build one in Phase 05; PyTorch's is the reference design).
5. Forgetting pinned memory for async: cudaMemcpyAsync from pageable memory quietly serializes.

3. Shared memory: the two-line contract

• Tile global→smem with coalesced cooperative loads; __syncthreads(); compute from smem; __syncthreads(); next tile. Matmul AI goes from O(1) to O(T): the whole game in one sentence.
• Pad [32][33] to dodge bank conflicts; never put __syncthreads() under a block-divergent branch.

4. Numbers worth carrying

5. The profiler triage (commit to memory)

nsys timeline:
  GPU idle?           → pipeline problem: launches, syncs, transfers. Stop kernel-tuning.
  GPU busy → ncu the top kernel:
    SOL Memory high?  → coalescing → reuse/tiling → vectorize loads (float4)
    SOL Compute high? → precision, Tensor Cores, better algorithm
    both low?         → latency-bound: occupancy, ILP, dependency chains

The strongest habit this phase teaches: predict before measuring. Write down the expected bandwidth/speedup from the roofline, then check. A team that predicts-then-profiles compounds; a team that profiles-then-rationalizes doesn't.

6. Debugging war stories (so you recognize them)

• "It crashes on a line that's obviously fine" — async error from an earlier kernel. CUDA_LAUNCH_BLOCKING=1, find the real culprit.
• "Works at n=1000, corrupts at n=1000000" — missing tail guard; or grid dim overflowed 65,535 in a .y/.z dimension.
• "Slower after I 'optimized' registers" — -maxrregcount caused spills; occupancy went up, performance went down. Volkov was right.
• "First call takes 200 ms" — context init + module JIT. Warm up; ship SASS for your target arch (Phase 03's fatbin story).
• "Results differ run to run by 1e-6" — float reduction order changes with block scheduling; that's not a bug, it's floating-point non-associativity. Define tolerance-based verification (and lock seeds for argmax-style outputs).
• "Everything returns errors after one bad kernel" — poisoned context; restart the process. Architect tenant isolation accordingly (Phase 06).

7. From engineer to head-of-engineering

What changes at your level isn't the syntax — it's the standards:

• Kernel PR checklist: reference-comparison test, sanitizer-clean, event-based timing with reps, ncu before/after naming the limiter, arch notes. (Lab 02 builds it; your job is making it stick.)
• Estimate sanity: when someone says "we'll get 10× by porting to GPU," you now ask: what's the AI? what fraction is parallel? how much PCIe traffic? Amdahl + roofline kills most 10× claims in five minutes — kindly, with math.
• Hiring signal: ask candidates the reduction-evolution question (§Ch. 10 of the WARMUP). Depth on "why shuffle beats shared memory for the last warp" separates practitioners from tutorial-readers.

Next: Phase 03 — what nvcc actually does to this code (PTX, SASS, and a mini compiler you'll write yourself), because a platform that promises "independence from underlying hardware" lives or dies in the compiler layer.