Lab 02 — Coalescing, Bank Conflicts & the Profiler

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

Concept primer: ../WARMUP.md Ch. 6–8, 11.

Run

make && ./bench               # bandwidth tables: stride, AoS/SoA, transpose ladder
ncu --set basic ./bench       # then: the profiler exercise (see §3)

0. The mission

Three experiments that turn WARMUP Chapter 6–7 claims into numbers you measured, then a profiler session that teaches you to name limiters with evidence:

  1. Stride sweep: copy bandwidth at stride 1, 2, 4, 8, 16, 32 — watch coalescing die.
  2. AoS vs SoA: load one field from a 3-float struct array vs a flat array.
  3. Transpose ladder: naive → shared-memory tiled → tiled+padded; recover copy bandwidth step by step.

Expected shape of results (T4; absolute numbers vary):

== stride sweep (effective GB/s) ==
stride  1: 240.1     stride  4:  68.3     stride 16:  17.9
stride  2: 124.5     stride  8:  34.6     stride 32:   9.2     <- ~26x cliff

== AoS vs SoA (loading .x only) ==
AoS:  81.4 GB/s   SoA: 238.0 GB/s   (~3x: the other 8 bytes ride the sectors)

== transpose ladder ==
copy (ceiling) : 235 GB/s
naive          :  62 GB/s   (writes uncoalesced)
tiled          : 148 GB/s   (both phases coalesced; bank conflicts on read)
tiled + pad    : 221 GB/s   (94% of copy — done)

1. Why each step behaves as it does

  • Stride s: each warp load touches ~min(32, s·4·32/32) distinct 32-byte sectors; useful bytes per sector fall as 1/s until one float per sector (s ≥ 8) — after which the cliff flattens: you're paying one full sector per element. The same curve you produced on the CPU in Phase 01 Lab 02, now at warp granularity.
  • AoS: reading pts[i].x drags .y and .z through the memory system as sector ballast. SoA puts 32 consecutive .x in 4 sectors. (This single diagram is why every GPU-resident data structure in Phases 05–07 is SoA.)
  • Transpose naive: reads coalesce (row-major in), writes scatter (column-major out) — one direction always loses. Tiled: stage a 32×32 tile in shared memory; read coalesced, write coalesced, the "turn" happens in smem. Padding [32][33]: the tiled version's smem reads hit one bank 32 ways; the +1 column shifts each row's bank phase → conflict-free. One character, ~1.5× on the transpose.

2. Record your tables

Fill in the actual numbers in PROFILE.md §1–3. Predict each before running — prediction-first is the habit.

3. The profiler session (§4 of PROFILE.md)

On the naive and padded transpose kernels:

ncu --kernel-name regex:transpose --set basic ./bench

Find and record, for each kernel:

Metric (section)What it tells you
Memory Throughput % / Compute (SM) % (Speed Of Light)which roof you're under
l1tex__average_t_sectors_per_request (Memory Workload)coalescing quality: ~4 good, ~32 catastrophic
shared_load_bank_conflicts (or l1tex__data_bank_conflicts...)the padding fix, visible
Achieved Occupancy (Occupancy)latency-hiding budget in use

Then write the verdict line for each kernel in PROFILE.md: "<kernel> is <memory/compute/latency>-bound because <two metrics>; the next experiment would be <X>." — that sentence format is the whole point of the lab. It's also the PR-review standard you'll set in Phase 12.

On Colab: !ncu --set basic ./bench works on T4 instances (driver permitting; if ERR_NVGPUCTRPERM, add --target-processes all or use nsys profile --stats=true ./bench for the timeline-level view instead and note the difference in what you can conclude).

4. Extensions

  • float4 vectorized copy: 16-byte loads per thread; measure against stride-1 float copy. Why does it help? (Fewer instructions per byte; wider sectors per request.)
  • cp.async (sm_80+): async global→shared copies for the tiled transpose; overlap the next tile's load with this tile's store.
  • Occupancy experiment: add __launch_bounds__(1024, 1) to the padded transpose, check -Xptxas=-v and achieved occupancy, explain the change.
  • Write the checklist: distill §3 into a one-page "kernel PR review checklist" — you will reuse it verbatim in Phase 12's leadership artifacts.

5. Common pitfalls

  1. Comparing kernels at different clocks (thermal throttle mid-run): lock with nvidia-smi -lgc where permitted, or interleave reps.
  2. Profiling with --set full on a long benchmark — replay overhead makes it crawl; use basic + targeted metrics.
  3. Concluding "memory-bound" from low compute% alone — check both SOL numbers; both-low means latency-bound, a different fix.
  4. Forgetting that ncu serializes kernels — overlap effects you saw in nsys vanish under ncu; the tools answer different questions (WARMUP Ch. 11).

6. What this lab proves about you

You can quantify the two memory sins (uncoalesced access, bank conflicts), fix them with the two standard remedies (SoA/tiling, padding), and defend the analysis with profiler evidence — the exact competency a GPU platform lead needs to review kernel work credibly.