Phase 02 — CUDA Programming: Kernels, Memory, Streams
Difficulty: ⭐⭐⭐⭐☆ | Estimated Time: 2 weeks Roles supported: Head of Engineering [GPU], GPU Platform Engineer, Kernel/Performance Engineer Hardware needed: any NVIDIA GPU — a free Google Colab T4 is enough for every lab here
Why This Phase Exists
The JD says "stay hands-on where needed — maintaining technical credibility with engineering teams." For a GPU platform, that credibility has one test: can you write, debug, and profile a CUDA kernel yourself? Not because the Head of Engineering writes kernels daily, but because every estimate you approve, every "it's a kernel problem" escalation you arbitrate, and every hire you evaluate depends on having done it.
Phase 01 gave you the machine model. This phase makes you operate the machine: the CUDA programming model (grids/blocks/threads), the memory API and its async variants, shared-memory tiling, coalescing, occupancy, streams, and the profiler that tells you the truth.
Concepts
- The CUDA software stack: runtime API vs driver API; what
nvccproduces (handoff to Phase 03) - Kernel launch anatomy: grid/block/thread,
<<<blocks, threads, smem, stream>>>, dimension math - The block-to-SM contract: why blocks must be independent; what that buys the hardware
- Error handling discipline:
cudaGetLastError, sync vs async errors, theCUDA_CHECKmacro - Memory management:
cudaMalloc/cudaMemcpy, pinned (page-locked) memory, unified memory and its page-fault cost - Global-memory coalescing: 32-byte sectors, AoS vs SoA, the strided-access cliff
- Shared memory:
__shared__, tiling, bank conflicts,__syncthreads()discipline - Occupancy in practice: registers/thread,
__launch_bounds__, the occupancy calculator - Streams and events: async copies, overlap,
cudaEventElapsedTimefor honest timing - Atomics and reductions: warp shuffle (
__shfl_down_sync), block reduction patterns - Profiling: Nsight Systems (timeline) vs Nsight Compute (kernel deep-dive); the three metrics that matter first
Labs
Lab 01 — First Kernels: Vector Add → Tiled Matmul (CUDA C++)
| Field | Value |
|---|---|
| Goal | Go from hello kernel to a shared-memory tiled matmul that beats the naive version ~5–10×, verified against CPU reference. |
| Concepts | Launch config math, error checking, coalescing, shared-memory tiling, __syncthreads(), honest benchmarking with events. |
| Steps | 1) Build & run make && ./kernels. 2) Read kernels.cu section by section against the README walkthrough. 3) Verify the naive→tiled speedup. 4) Do the extension ladder (rectangular tiles, register blocking, compare to cuBLAS). |
| Stack | CUDA C++ (nvcc), works on any sm_70+ GPU; Colab notebook instructions included |
| Output | Working binary + a timing table: naive vs tiled vs (extension) cuBLAS. |
| How to Test | Built-in: every kernel's output is compared elementwise against a CPU reference; speedup asserts print PASS/FAIL. |
| Talking Points | Why the tiled version is faster (AI raised from ~O(1) to ~tile-size — the roofline from Phase 01 Lab 02); why __syncthreads() twice per tile; what cuBLAS still does that you don't. |
| Resume Bullet | "Implemented and verified CUDA kernels from vector-add through shared-memory tiled GEMM; achieved 8× speedup over naive matmul and validated the roofline prediction with Nsight Compute metrics." |
| Extensions | Register blocking (each thread computes a 2×2 output micro-tile); compare against cuBLAS sgemm and report the remaining gap; try __launch_bounds__ and explain the occupancy change. |
Lab 02 — Coalescing, Occupancy & the Profiler (CUDA C++)
| Field | Value |
|---|---|
| Goal | Measure the cost of uncoalesced access and bank conflicts; read Nsight Compute output like a platform engineer. |
| Concepts | 32-byte sectors, strided/transposed access, AoS vs SoA, shared-memory banks + padding fix, occupancy limiters. |
| Steps | 1) make && ./bench. 2) Record the bandwidth table (stride 1/2/8/32; AoS vs SoA; transpose naive vs tiled vs padded). 3) Profile two kernels with ncu and identify the limiter (memory vs compute vs latency). 4) Fill in PROFILE.md. |
| Stack | CUDA C++, Nsight Compute CLI (ncu) — Colab-compatible commands provided |
| Output | Bandwidth/conflict tables + a filled PROFILE.md naming each kernel's limiter with evidence. |
| How to Test | Stride-32 bandwidth collapses to ~1/8–1/32 of stride-1; padded transpose recovers ≥80% of copy bandwidth. |
| Talking Points | "Coalescing is the GPU's cache-line story" (ties to Phase 01 Lab 02 strided extension); why SoA wins on GPU; how you'd review a kernel PR with only an ncu report attached. |
| Resume Bullet | "Quantified memory-coalescing and bank-conflict penalties on Ampere (8–20× bandwidth loss); established Nsight-Compute-based review checklist adopted for kernel PRs." |
| Extensions | Vectorized loads (float4) — measure the gain; async copy (cp.async, sm_80+) into shared memory; build a one-page "kernel review checklist" for your team. |
Deliverables Checklist
- Both labs build and pass their built-in verification on a real GPU
- Timing table: naive vs tiled matmul, with the roofline explanation written down
- Bandwidth tables for stride / AoS-SoA / transpose experiments
-
PROFILE.mdfilled: each kernel's limiter named withncuevidence -
You can write the
CUDA_CHECKmacro and the launch-config formula from memory
Interview Relevance
- "Launch config for a 1920×1080 image filter — walk me through the math."
- "Your kernel gets 8% of peak bandwidth. List your first three hypotheses."
- "Why are thread blocks required to be independent? What feature does that enable?"
- "What does
__syncthreads()do, and what happens if it's inside a divergent branch?" - "When is unified memory acceptable in production?"