reduction, shuffle, atomics, numerical accuracy
Parallel Reduction: Warp, Block, and Grid Levels
REFERENCE CONTEXT · CUDA 12.2; RTX 2070 SUPER; Turing sm_75
Core Concept
Efficient reductions are hierarchical: local thread accumulation, warp shuffle, one shared value per warp, then a final block reduction. Grid-wide reduction usually needs a second kernel or atomics.
Why It Affects Performance
This structure avoids repeated block-wide barriers and unnecessary shared-memory traffic while retaining parallelism.
Example or Pseudocode
float sum = 0.f;
for (int i = global_tid; i < n; i += grid_stride) sum += x[i];
sum = warp_reduce(sum);
if ((threadIdx.x & 31) == 0) warp_sums[threadIdx.x >> 5] = sum;Profiling and Performance Judgment
Measure GB/s, barrier stalls, atomic throughput, and error. Test non-power-of-two sizes because floating-point reassociation changes results.
Five-Minute Exercise
How many warp partials does a 256-thread block produce?
Show answer
Eight. The final stage must reduce eight warp sums.