Lab 01 — First Kernels: Vector Add → Tiled Matmul

Phase: 02 — CUDA Programming | Difficulty: ⭐⭐⭐☆☆ | Time: 4–6 hours Language: CUDA C++ | Hardware: any NVIDIA GPU (sm_70+); Colab T4 works

Concept primer: ../WARMUP.md Ch. 2–4, 7, 9.

Run

make            # nvcc -O3 -arch=native (or ARCH=sm_75 make for Colab T4)
./kernels

No local GPU? Google Colab

Runtime → Change runtime type → T4 GPU, then in one cell:

!nvcc --version
%%writefile kernels.cu
<paste kernels.cu>
!nvcc -O3 -arch=sm_75 -o kernels kernels.cu && ./kernels

0. The mission

Three kernels, one discipline:

  1. vec_add — the programming model + error handling + the tail guard.
  2. matmul_naive — one thread per output element, all reads from global memory.
  3. matmul_tiled — the same math with shared-memory tiling.

Every kernel is verified elementwise against a CPU reference, and timed with CUDA events (warmup + median of reps). Expected result: tiled beats naive ~5–10×, because tiling raises arithmetic intensity by the tile factor — the Phase 01 roofline argument made executable.

Sample output (T4, n=1024):

vec_add        : PASS   33.55 GB/s effective (memory-bound, as designed)
matmul_naive   : PASS   612 ms   0.35 TFLOP/s
matmul_tiled   : PASS    78 ms   2.75 TFLOP/s   speedup 7.8x  [PASS >= 4x]

1. Read the code in this order

  1. CUDA_CHECK and check_close — the discipline. Note where cudaGetLastError is called (right after launch — catches config errors) vs cudaDeviceSynchronize (catches the kernel's own faults).
  2. vec_add — the index formula and tail guard. Exercise: delete the if (i < n) guard, run with n=1<<20+3, observe either corruption or a sanitizer report (compute-sanitizer ./kernels). Put it back.
  3. matmul_naive — each thread reads a row of A and a column of B from global memory. The column read of B is actually coalesced across threads (adjacent tx → adjacent B columns at fixed k) — the problem isn't coalescing here, it's zero reuse: every element fetched N times. AI ≈ O(1).
  4. matmul_tiled — the cooperative load / sync / compute / sync rhythm (WARMUP Ch. 7). Exercise: comment out the second __syncthreads() and run — verification fails intermittently (a fast warp overwrites the tile while a slow one still reads). This is the cheapest race-condition education available; take it.
  5. The timing harness — events, warmup, median-of-5. Compare against the naive std::chrono around an async launch (also included, prints the lie).

2. Why tiled wins — write this down in your own words

Naive: each of N³ multiply-adds needs 2 fresh global loads → AI ≈ 1/8 FLOP/byte → the bandwidth roof at ~320 GB/s (T4) caps you near 0.04 TFLOP/s per... except L2 catches some reuse, so you observe ~0.3. Tiled with T=32: each global element is loaded once per tile and reused 32 times from shared memory → AI × 32 → the roof rises past 1 TFLOP/s on the same silicon. Same FLOPs, same GPU, 8× — layout and staging, not arithmetic.

The remaining 3–5× to cuBLAS: register micro-tiles (each thread computes 4–8 outputs from registers, raising AI again), double-buffered async tile loads (cp.async), Tensor Cores via mma, and per-arch tile-shape tuning. That ladder is CUTLASS's table of contents — and the extension exercises walk its first rung.

3. Extension ladder

  1. Rectangular blocking: 64×16 tiles; measure. Why might non-square help? (Hint: B's tile loads per output column.)
  2. Register micro-tile: each thread computes a 2×2 patch of C. Expect another 1.5–2.5×. You're now register-blocked like real GEMMs.
  3. cuBLAS comparison: link -lcublas, time cublasSgemm, report your % of it. >30% with micro-tiles is respectable; explain the rest.
  4. __launch_bounds__(256, 4) on the tiled kernel: check -Xptxas=-v register count before/after, occupancy, and time. Did the compiler spill?

4. Common pitfalls

  1. Host pointer passed to a kernel — illegal access at the next sync, blamed on the wrong line (CUDA_LAUNCH_BLOCKING=1 to localize).
  2. Verifying with == on floats — reduction/FMA order differs from CPU; use relative tolerance (the harness uses 1e-3 relative).
  3. Timing the enqueue, not the kernel (chrono without sync) — the harness prints both so you see the discrepancy once and never trust chrono again.
  4. -arch mismatch — sm_90 binary on a T4 dies at load; -arch=native locally, explicit ARCH= on shared machines (the full fatbin story is Phase 03).

5. What this lab proves about you

You can produce a verified, honestly-benchmarked kernel and explain its performance from the roofline — the bar you will later hold every kernel PR to (Lab 02 turns that into a checklist).