Warmup Guide — CUDA Programming

Zero-to-expert primer for Phase 02. Assumes Phase 01's machine model (warps, SMs, memory hierarchy, roofline) and basic C/C++. By the end you can write, verify, time, and profile real kernels — and review other people's.

Table of Contents

Chapter 1: The CUDA Stack — What You're Actually Calling

Zero background: "CUDA" is three things people conflate: a language extension (C++ plus __global__, <<<>>>), a runtime library (libcudart — the cudaMalloc/cudaMemcpy API), and a driver stack (libcuda.so, the kernel-mode driver — Phase 04 territory).

The call chain when you launch a kernel:

your code → CUDA Runtime API (libcudart)   cudaLaunchKernel(...)
          → CUDA Driver API   (libcuda.so) cuLaunchKernel(...)
          → ioctl(2) into the kernel-mode driver (/dev/nvidia*)
          → commands written to a ring buffer the GPU hardware consumes

Two facts with production consequences:

  1. The runtime API is a convenience layer over the driver API. Inference engines and runtimes (Phase 05/09) often use the driver API directly for explicit context control, loading cubins at runtime (cuModuleLoad), and multi-tenancy. You should be able to read both.
  2. Kernel launches are asynchronous. The CPU enqueues and returns in ~5–10 µs; the GPU executes later. Everything about timing (Ch. 9), error reporting (Ch. 4), and overlap (streams) follows from this one fact.

Misconception: "nvcc is a compiler like gcc." nvcc is a driver that splits your file: host code → your host compiler; device code → PTX/SASS through NVIDIA's backend. Phase 03 dissects this pipeline; here you just use it.

Chapter 2: The Programming Model — Grid, Block, Thread

You write one scalar function (a kernel); the runtime launches N instances:

  • Thread: one instance. Knows its coordinates via threadIdx, blockIdx, blockDim, gridDim (each a 3-component vector for 1D/2D/3D convenience).
  • Block: up to 1,024 threads that share a shared-memory allocation, can __syncthreads(), and are guaranteed to run on one SM.
  • Grid: all blocks of a launch. Blocks must be independent — no cross-block sync within a kernel, no assumption about execution order.

Why block independence is the load-bearing rule: it's what lets the hardware scheduler scatter blocks across however many SMs exist — 10 on a laptop, 132 on H100 — and retire/replace them in any order. Block independence is CUDA's scalability contract; it is why 2008 kernels still scale on 2025 silicon. (It's also why a "wait for block X" spinlock in a kernel deadlocks: block X may not be resident.) Cooperative groups (cudaLaunchCooperativeKernel) exist as the explicit, occupancy-checked exception.

The index formula you'll write a thousand times:

int i = blockIdx.x * blockDim.x + threadIdx.x;     // global index
if (i < n) { ... }                                  // guard the tail
int nblocks = (n + threads - 1) / threads;          // ceil-div launch config

The guard exists because the grid is a multiple of the block size and n usually isn't. Forgetting it = out-of-bounds writes = the classic first CUDA bug (intermittent corruption, not a crash — see Ch. 4).

Mapping back to Phase 01: blocks are carved into warps (threads 0–31, 32–63, …). All divergence/coalescing reasoning operates at warp granularity within the block you configured. Block size should essentially always be a multiple of 32 — 128 or 256 are the standard defaults.

Chapter 3: Your First Kernel, Line by Line

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}
// host:
int threads = 256;
int blocks  = (n + threads - 1) / threads;
vec_add<<<blocks, threads>>>(d_a, d_b, d_c, n);
  • __global__: callable from host, runs on device. (__device__ = device-only helper; __host__ __device__ = compiled for both.)
  • The pointers are device pointers (from cudaMalloc). Dereferencing a host pointer in a kernel is undefined behavior — typically illegal memory access reported much later (Ch. 4).
  • <<<blocks, threads>>> has two more slots: dynamic shared-memory bytes and a stream — <<<blocks, threads, smem_bytes, stream>>> (Ch. 7, 9).

Production reality: vector add is the memory-bound kernel (AI ≈ 0.08, Phase 01 Ch. 7). Its only performance question is "did we hit memory bandwidth?" — which makes it the perfect coalescing test harness, and exactly how Lab 01 uses it.

Chapter 4: Error Handling — The Discipline That Separates Professionals

CUDA errors come in two flavors, and conflating them wastes engineer-days:

  • Synchronous: bad launch config, OOM at cudaMalloc — returned immediately by the call.
  • Asynchronous: the kernel itself faults (out-of-bounds, bad pointer) — surfaces at the next synchronizing call, attributed to an innocent line.

The non-negotiable macro (write it from memory in interviews):

#define CUDA_CHECK(call) do {                                         \
    cudaError_t e = (call);                                           \
    if (e != cudaSuccess) {                                           \
        fprintf(stderr, "CUDA error %s:%d: %s\n",                     \
                __FILE__, __LINE__, cudaGetErrorString(e));           \
        exit(1);                                                      \
    }                                                                 \
} while (0)

kernel<<<g, b>>>(args);
CUDA_CHECK(cudaGetLastError());        // catches launch-config errors
CUDA_CHECK(cudaDeviceSynchronize());   // catches kernel-execution errors (debug builds)

Debugging ladder when something faults: 1) compute-sanitizer ./app (catches the faulting kernel + address — the GPU's valgrind), 2) CUDA_LAUNCH_BLOCKING=1 (makes async errors synchronous so attribution is correct), 3) cuda-gdb for the rare hard case.

Production significance: a fault in any kernel poisons the CUDA context — subsequent calls return errors until process restart. This is why serving engines treat "CUDA error" as fatal-and-restart (Phase 10 runbooks), and why your platform must isolate tenants at process granularity, not context granularity (Phase 06).

Chapter 5: The Memory API — cudaMalloc to Unified Memory

The four tiers, in order of how often you should reach for them:

  1. Explicit — cudaMalloc + cudaMemcpy. Predictable, fast, verbose. The production default. cudaMemcpy is synchronous-ish; the async variant requires tier 2.
  2. Pinned host memory — cudaMallocHost. Page-locked so the GPU's DMA engine can read it directly: ~2× faster transfers and the prerequisite for cudaMemcpyAsync actually being async (pageable copies silently stage through a pinned bounce buffer and serialize). Cost: pinned pages are stolen from the OS; pin gigabytes, not tens of gigabytes.
  3. Unified memory — cudaMallocManaged. One pointer valid on both sides; pages migrate on demand. Beautiful for prototypes; in production the page-fault storms on first touch and on every CPU↔GPU ping-pong are a classic mystery slowdown (visible as Unified Memory rows in Nsight Systems). cudaMemPrefetchAsync is the remedy when you must use it.
  4. Stream-ordered pools — cudaMallocAsync/cudaFreeAsync. Allocation as a stream operation from a driver-managed pool — because raw cudaMalloc/Free synchronize the device and fragment, which is also why every framework built its own caching allocator. Phase 05 Lab 01 has you build exactly that allocator in Rust — this chapter is its motivation.

The rule that survives all four tiers: minimize transfers; batch them; overlap them with compute (Ch. 9). PCIe is 50× narrower than HBM (Phase 01 Ch. 6) — a synchronous host↔device copy inside a hot loop is the single most common production CUDA performance bug.

Chapter 6: Coalescing — The #1 Performance Rule

The mechanism: when a warp executes a load, the hardware takes all 32 addresses and computes which 32-byte sectors of memory they touch, then issues one transaction per distinct sector.

  • 32 consecutive floats → 4 sectors → minimal transactions. Coalesced.
  • Stride-32 floats (e.g., walking a column of a row-major matrix) → 32 different sectors → 32 transactions, each delivering 4 useful bytes of 32 fetched. ~8× wasted bandwidth, and worse with bigger strides.

This is the GPU's version of Phase 01 Lab 02's cache-line stride cliff — same physics, warp-granular.

The three patterns to know:

  1. Adjacent threads ↔ adjacent addresses. a[blockIdx.x*blockDim.x + threadIdx.x] is perfect; a[threadIdx.x * pitch] is the anti-pattern.
  2. SoA beats AoS. struct {float x,y,z;} pts[n] makes pts[i].x loads stride-12; three arrays x[n], y[n], z[n] coalesce each field. This is why GPU-facing data layouts (and Phase 07's KV-cache layouts) are structure-of-arrays.
  3. Matrix transpose is the canonical exercise: naive transpose must uncoalesce either reads or writes; the fix is staging a tile through shared memory so both global phases are coalesced — Lab 02 implements it, padding included (Ch. 7).

Misconception: "the L2 cache makes coalescing obsolete." L2 softens repeated sloppiness, but a bandwidth-bound kernel wasting 8× on sectors is still ~8× slow — Lab 02's table is the proof you'll carry around.

Chapter 7: Shared Memory and Tiling — Raising Arithmetic Intensity

Phase 01 Ch. 7 said: performance = min(peak, AI × bandwidth). Tiling is the technique for raising AI, and shared memory is its vehicle.

Naive matmul: each thread computes one C[i][j], reading a full row of A and column of B from global memory → every input element fetched N times → AI ≈ O(1) → bandwidth-bound at a few % of peak FLOPs.

Tiled matmul (Lab 01's centerpiece):

__shared__ float As[T][T], Bs[T][T];
for (int t = 0; t < N/T; t++) {
    As[ty][tx] = A[row*N + t*T + tx];      // coalesced cooperative load
    Bs[ty][tx] = B[(t*T + ty)*N + col];    // coalesced cooperative load
    __syncthreads();                        // tile fully loaded before use
    for (int k = 0; k < T; k++) acc += As[ty][k] * Bs[k][tx];
    __syncthreads();                        // all done reading before overwrite
}

Each global element is now loaded once per tile instead of once per thread → AI multiplied by T (tile dim, typically 16–32) → 5–10× speedup, exactly as the roofline predicts. cuBLAS continues the same idea further down the hierarchy: register-blocking (each thread computes a small output micro-tile from registers) and Tensor Core fragments — that's the remaining gap when Lab 01's extension compares against sgemm.

The two rules of __syncthreads(): (1) it's a block-wide barrier — every thread of the block must reach it, so it must never sit inside a branch that's divergent across the block (deadlock/UB); (2) you need it twice per tile — after loading (before consuming) and after consuming (before the next overwrite). Lab 01 asks you to remove one and observe the corruption.

Bank conflicts: shared memory has 32 banks, word-interleaved. tile[threadIdx.x][k] with a 32-wide tile puts a warp's accesses in one bank (32-way conflict → serialized). Fix: pad the tile — __shared__ float tile[32][33] — the +1 rotates each row's bank alignment. Lab 02 measures before/after; it's a ~30–40% transpose-bandwidth swing for one character of code.

Chapter 8: Occupancy in Practice

Phase 01 Ch. 5 defined occupancy as a latency-hiding budget. The practice:

What limits it (per SM, Ampere numbers): 65,536 registers, 100–164 KB shared memory, 2,048 thread slots, 32 block slots. Whichever resource your kernel exhausts first caps how many blocks are resident:

  • 256 threads/block, 64 regs/thread → 16,384 regs/block → 4 blocks → 1,024 threads = 50% occupancy (register-limited).
  • 48 KB smem/block on a 100 KB SM → 2 blocks resident regardless of registers.

The tools: nvcc -Xptxas=-v prints regs/smem per kernel; cudaOccupancyMaxActiveBlocksPerMultiprocessor computes it at runtime (use it to pick launch configs in library code — Phase 09's HAL does this); __launch_bounds__(maxThreads, minBlocks) tells the compiler to fit a target (it will spill registers to comply — measure, don't assume).

The judgment call (this is what reviewers actually argue about): more occupancy hides more latency, but fewer registers per thread can mean spills, and less shared memory per block can mean less tiling/reuse. There is no universally right answer — only the profile. Volkov again: high-ILP kernels can peak at 25% occupancy. The professional posture: state the limiter, show the ncu evidence, justify the chosen point.

Chapter 9: Streams, Events, and Honest Timing

Streams are ordered work queues; operations in different streams may overlap. The legacy default stream synchronizes with everything (compile with --default-stream per-thread or just use explicit streams everywhere — the platform-code default).

The canonical overlap pattern — pipelining chunked transfers against compute (this is also Phase 07's prefill/decode overlap in miniature):

for (int c = 0; c < nchunks; c++) {
    cudaMemcpyAsync(d_in + off, h_in + off, bytes, H2D, stream[c % 2]);
    kernel<<<g, b, 0, stream[c % 2]>>>(d_in + off, d_out + off);
    cudaMemcpyAsync(h_out + off, d_out + off, bytes, D2H, stream[c % 2]);
}

Requirements for actual overlap: pinned host memory (Ch. 5), separate streams, and chunk sizes big enough to amortize launch overhead (~10 µs/launch — relevant to Phase 05's scheduler and why CUDA Graphs exist for short-kernel pipelines).

Honest timing — the three rules, violated in 90% of bad benchmark claims:

  1. Time with events (cudaEventRecord/cudaEventElapsedTime), or wall-clock only around an explicit cudaDeviceSynchronize(). Timing an async launch with std::chrono measures the enqueue, not the kernel.
  2. Warm up first (first launch includes module load/JIT — Phase 03 explains the JIT path).
  3. Report the median of many reps, with clocks noted if comparing across sessions (GPUs thermal-throttle; nvidia-smi -q -d CLOCK).

Events also serve as cross-stream dependencies (cudaStreamWaitEvent) — the primitive from which Phase 05 Lab 02's stream scheduler builds its DAG.

Chapter 10: Reductions and Warp Primitives

Summing N elements exposes everything this phase taught at once — and it's a top-3 interview exercise.

The evolution every CUDA engineer should be able to narrate:

  1. Atomic per element: atomicAdd(&out, a[i]) — correct, serializes on one address, terrible.
  2. Shared-memory tree: block loads to smem, halving strides with __syncthreads() between levels, thread 0 atomically adds the block's sum. Classic; bank-conflict-aware versions use sequential addressing (stride from blockDim/2 down, not interleaved by 2× — the interleaved version also has divergence problems: if (tid % (2*s) == 0) is the per-lane branch Phase 01 warned about).
  3. Warp shuffle finish: the last 32 elements don't need shared memory at all — val += __shfl_down_sync(0xffffffff, val, offset) moves data register-to-register within the warp, no smem, no sync. Modern reductions: shuffle within warps, one smem round between warps, shuffle again.
  4. (Library answer: CUB BlockReduce — and saying "in production I'd use CUB, here's what it does inside" is the strongest possible interview answer.)

Warp primitives to know by name: __shfl_*_sync (data exchange), __ballot_sync / __any_sync / __all_sync (votes — Lab 01 of Phase 01 simulated vote.any), __activemask(). All take the mask-first argument because of Volta independent thread scheduling (Phase 01 Ch. 3).

Chapter 11: Profiling — Nsight Systems vs Nsight Compute

Two tools, two altitudes — using the wrong one wastes afternoons:

  • Nsight Systems (nsys) — the timeline. CPU threads, API calls, transfers, kernel spans, stream overlap. Answers: "where does wall-clock time go?" "is the GPU idle between kernels?" "did my copies overlap?" Always start here. nsys profile -o report ./app → open in the GUI or nsys stats.
  • Nsight Compute (ncu) — the microscope. Per-kernel hardware counters: achieved bandwidth, warp execution efficiency (Lab 01 of Phase 01 computed this by hand!), occupancy achieved vs theoretical, stall reasons, and a built-in roofline chart. ncu --set full ./app. Expensive (replays kernels many times).

The first-pass triage (memorize as a flowchart):

  1. nsys: GPU mostly idle? → it's a pipeline problem (launch overhead, sync points, transfers) — no kernel tuning will help.
  2. GPU busy → ncu the top kernel: SOL Memory% vs SOL Compute% tells you which roof you're under (Phase 01's roofline, now measured for you).
  3. Memory-bound → check gld_efficiency/sectors-per-request (coalescing, Ch. 6), then data reuse (tiling, Ch. 7). Compute-bound → precision/Tensor-Core utilization. Neither high but stalls high → latency-bound: occupancy (Ch. 8) or dependency chains.

This triage is Lab 02's PROFILE.md exercise, and — at the leadership level — it's the checklist you require attached to any "we optimized the kernel" PR.

Lab Walkthrough Guidance

Order: Lab 01 → Lab 02. Lab 01 builds the kernels; Lab 02 measures their sins.

Lab 01 (first kernels):

  1. Build and run first (make && ./kernels); confirm all PASS lines, then read kernels.cu against Chapters 2–4 and 7.
  2. Work the verification path: break the tail guard deliberately, see the corruption, fix it, then break a __syncthreads() and see the race.
  3. Extensions in order: rectangular tiles → register micro-tiles → cuBLAS comparison (expect cuBLAS 2–5× ahead; explain why using Ch. 7's last paragraph).
  4. No local GPU: the README's Colab cell runs everything; nvcc flags identical.

Lab 02 (coalescing & profiler):

  1. Run the benchmark and fill the tables before profiling — predict, then verify with ncu (prediction-first is the habit that makes profiling stick).
  2. The transpose ladder (naive → tiled → padded) is the heart: tie each step to Ch. 6/7 mechanisms in one sentence each in PROFILE.md.
  3. Finish with the triage flowchart (Ch. 11) applied to both your matmuls from Lab 01 — which roof is each under, with evidence.

Success Criteria

  • You can write a correct kernel + launch + CUDA_CHECK discipline from a blank editor with no references (Ch. 2–4)
  • You can explain sync vs async errors and run the 3-step debugging ladder (Ch. 4)
  • You can state the coalescing mechanism in terms of warp addresses → 32-byte sectors, and the SoA and transpose consequences (Ch. 6)
  • Your tiled matmul beats naive ~5–10× and you can derive why from the roofline, including both __syncthreads() placements (Ch. 7)
  • You can name a kernel's occupancy limiter from -Xptxas=-v output (Ch. 8)
  • Your benchmark numbers come from events with warmup and reps (Ch. 9)
  • You can narrate the reduction evolution through warp shuffles (Ch. 10)
  • Given an nsys+ncu pair, you can name the limiter and the next experiment in under two minutes (Ch. 11)

Interview Q&A

Q1: Walk me through launching a blur filter over a 1920×1080 image. A: Pick a 2D block, say 16×16 = 256 threads. Grid = ceil(1920/16) × ceil(1080/16) = 120 × 68 blocks. Each thread computes x = blockIdx.x*16 + threadIdx.x, y likewise, guards if (x < 1920 && y < 1080) because 1080/16 = 67.5 rounds up and the last block row hangs off the edge. For a blur, neighboring threads read overlapping pixels, so I'd stage a (16+2r)² halo tile in shared memory with cooperative coalesced loads, __syncthreads(), then filter from smem. Launch asynchronously into a stream, CUDA_CHECK(cudaGetLastError()) after.

Q2: Your kernel achieves 8% of peak bandwidth. First three hypotheses? A: (1) Uncoalesced access — warp addresses scattering across 32-byte sectors: check ncu sectors-per-request (~4 is ideal for floats; 32 means fully strided) — typical causes: column-major walk, AoS layout, bad pitch. (2) Latency-bound, not bandwidth-bound — too few warps in flight to saturate the memory system: check achieved occupancy and stall reasons; fix with more parallelism per launch or more ILP per thread. (3) The kernel isn't actually the time sink — short kernel dominated by launch overhead or serialized with transfers: check the nsys timeline first. (Bonus: unified-memory page faults masquerading as kernel time.)

Q3: Why must thread blocks be independent? A: Independence is the scalability contract: the hardware can place, execute, and retire blocks in any order on any SM, which lets one binary saturate a 10-SM laptop and a 132-SM H100, and lets the scheduler backfill SMs as blocks finish (Phase 01's latency-hiding at block granularity). Inter-block sync within a kernel would require all blocks resident simultaneously — deadlock otherwise. The sanctioned escape hatches: end the kernel (grid-wide barrier = kernel boundary), atomics for unordered communication, or cooperative launch, which checks co-residency at launch time.

Q4: What goes wrong with __syncthreads() in divergent code? A: It's a block-wide barrier; every thread must execute the same barrier. If some threads of a block skip it (a branch divergent at block scope), the arriving threads wait forever — deadlock or undefined behavior, in practice a hang or corrupted results. Warp-level divergence reconverges by hardware, but __syncthreads() semantics demand all-threads participation: the rule is no __syncthreads() under conditions that vary within the block; restructure so the barrier is unconditional (compute the condition, sync, then branch).

Q5: Pinned vs pageable vs unified memory — when each? A: Pageable (malloc) host memory forces the driver to stage transfers through an internal pinned bounce buffer — slower, and cudaMemcpyAsync silently degrades to effectively synchronous. Pinned (cudaMallocHost) gives DMA-direct ~2× transfer speed and true async — required for any overlap pipeline; cost is stealing unpageable RAM from the OS, so pin transfer buffers, not datasets. Unified (cudaMallocManaged) trades control for convenience: on-demand page migration is fine for prototypes and oversubscription-with-prefetch on data too big for HBM, but fault storms on alternating CPU/GPU touches are a notorious production slowdown — in serving infrastructure I require explicit transfers at review time unless there's a measured case.

Q6 (leadership): An engineer's PR claims a 3× kernel speedup. What do you require before merge? A: (1) Correctness evidence: elementwise comparison against a reference within a stated tolerance, plus a compute-sanitizer clean run — fast-but-wrong kernels are the classic regression. (2) Honest methodology: event-based timing, warmup, median of N, fixed clocks noted, same input distribution as production. (3) The mechanism: which limiter moved, with before/after ncu evidence (e.g., sectors-per-request 32→4, SOL Memory 85%) — "I don't merge speedups we can't explain, because unexplained wins are usually measurement bugs." (4) Portability notes: which architectures it was tuned on, occupancy impact at other block sizes (Phase 09's HAL concern). This checklist is Lab 02's PROFILE.md, generalized into team process.

References

  • NVIDIA, CUDA C++ Programming Guide — https://docs.nvidia.com/cuda/cuda-c-programming-guide/
  • NVIDIA, CUDA C++ Best Practices Guide — https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
  • Mark Harris, "Optimizing Parallel Reduction in CUDA" (the classic 7-step deck) — https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
  • Mark Harris, "An Efficient Matrix Transpose in CUDA C/C++" — https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/
  • NVIDIA Nsight Systems / Nsight Compute documentation — https://docs.nvidia.com/nsight-systems/ , https://docs.nvidia.com/nsight-compute/
  • CUTLASS (how production GEMMs are actually structured) — https://github.com/NVIDIA/cutlass
  • Sanders & Kandrot, CUDA by Example (dated but the gentlest on-ramp)
  • Cheng, Grossman, McKercher, Professional CUDA C Programming
  • Cross-track: Phase 01 WARMUP (machine model), Phase 03 WARMUP (what nvcc does with this code)