LESSON 03 · 2026.04.18 · T4
Reduction — why atomic is 100× slower
Five implementations for sum(x). Shared memory, warp shuffle, and the 29% one __syncthreads made.
Five versions
- v1 atomic — every thread does
atomicAddon a single global address (baseline, bad example) - v2 shared — per-block shared-memory tree, only the root does an atomic
- v3 unroll — v2 + final warp replaced with shuffles, 5
__syncthreadsremoved - v4 shuffle — warps reduce locally via shuffle, shared memory only between warps
- thrust —
thrust::reduce, cub-based multi-pass
best_ms by version
| n | v1 | v2 | v3 | v4 | thrust |
|---|---|---|---|---|---|
| 2²⁰ | 2.082 | 0.018 | 0.013 | 0.012 | 0.028 |
| 2²² | 8.309 | 0.070 | 0.066 | 0.066 | 0.089 |
| 2²⁴ | 33.227 | 0.260 | 0.258 | 0.258 | 0.289 |
| 2²⁶ | 132.903 | 1.089 | 1.087 | 1.090 | 1.118 |
| 2²⁸ | 531.577 | 4.860 | 4.870 | 4.887 | 4.648 |
Lesson 1 · v1's "2 GB/s floor"
v1's effective bandwidth is pinned at ~2.0 GB/s regardless of n. That's not an HBM ceiling — it's an atomic throughput ceiling. When a million threads hit the same address simultaneously, the hardware serializes — effectively one thread.
Never use atomic as "the thing that produces one result." Use it once per block (the tree-reduction root), or as a low-frequency counter.
Lesson 2 · contribution of each step v2 → v3 → v4
| step | Δ time | speedup | cost removed |
|---|---|---|---|
| v1 → v2 | 2.082 → 0.018 ms | 113× | serial atomic → parallel tree |
| v2 → v3 | 0.018 → 0.013 ms | −29% | 5 __syncthreads |
| v3 → v4 | 0.013 → 0.012 ms | −6% | smem → register shuffle |
A single __syncthreads() costs hundreds of cycles. Five of them take up 1/3 of a small kernel's runtime. A number worth remembering.
Lesson 3 · at large n, everything converges
From n ≥ 2²⁴, v2/v3/v4/thrust all tie within ±5%. The reason is simple. HBM bandwidth is bottomed out (~77%) — once the bytes-to-read is fixed, "fetch time" dominates everything. No reduction tree layout shrinks that time.
Only worry about tail optimization at small n. At large n, prefer "code simplicity > in-kernel micro-optimization."
When does Thrust win?
Thrust is cub-based multi-pass: pass 1 writes per-block partial sums to a temp buffer, pass 2 finalizes. Overhead hurts at small n, HBM scheduling helps at large n — hence the crossover. Thrust always wins on accuracy — balanced tree depth cancels FP error more favorably.
Why write these by hand?
If the production answer is thrust/cub, the reason to hand-write v1–v4 is one thing: the gap in judgment between an engineer who remembers "why atomic is 100× slower" in numbers and one who doesn't. You need a reference point to evaluate Triton, Mojo, or the next compiler. Fusion decisions, bottleneck estimates, roofline reading — all of it runs on this sense.