Parallel Reduction Algorithm

Parallel Reduction is the fundamental parallel algorithm pattern that combines N input values into a single result (sum, max, min, dot product) using a logarithmic-depth tree of binary operations — requiring only O(log N) steps on N processors compared to O(N) for sequential reduction, making it the building block for virtually every aggregate computation in parallel and GPU computing.

The Reduction Tree

Given N = 1024 values to sum:
- Step 1: 512 threads each add pairs of adjacent elements → 512 partial sums
- Step 2: 256 threads add pairs of those partial sums → 256 values
- Step 3-10: Continue halving until 1 final sum remains
- Total: log2(1024) = 10 steps, with decreasing parallelism at each level

GPU Implementation Hierarchy

1. Warp-Level Reduction (Shuffle Instructions): Within a single warp (32 threads), CUDA's `__shfl_down_sync()` instruction allows threads to directly exchange register values without going through shared memory. A 32-element reduction completes in 5 shuffle operations (~5 cycles). This is the fastest reduction primitive.

2. Block-Level Reduction (Shared Memory): For thread blocks with multiple warps (e.g., 256 threads = 8 warps), each warp first reduces its 32 elements via shuffles, then the 8 warp results are combined via shared memory. The final value is written to global memory by one thread.

3. Grid-Level Reduction (Kernel Launch or Atomics): Multiple thread blocks each produce local sums. A second kernel launch (or atomic operations) combines the per-block results. Two-pass reduction (large kernel → small kernel) is standard for arrays with millions of elements.

Optimization Techniques

• Sequential Addressing (Avoid Divergence): Use stride = blockDim/2, blockDim/4, ... instead of stride = 1, 2, 4, ... to keep active threads in contiguous positions, avoiding warp divergence.
• First-Add During Load: Each thread loads and adds two (or more) elements during the initial global memory read, halving the number of threads needed and doubling memory throughput per thread.
• Unroll the Last Warp: When the number of active threads drops to 32, switch from shared-memory reduction to warp shuffles, eliminating synchronization overhead.
• Grid-Stride Loop: Each thread processes multiple elements in a loop before the tree reduction begins, maximizing work per thread and minimizing the number of thread blocks.

Performance

An optimized GPU reduction on an A100 achieves 80-90% of peak memory bandwidth (1.5+ TB/s) for large arrays — the operation is entirely memory-bandwidth-bound since the arithmetic (addition) is trivially cheap compared to the data movement.

Parallel Reduction is the simplest algorithm that exposes the full complexity of GPU optimization — just adding numbers reveals every performance pitfall: memory coalescing, warp divergence, shared memory bank conflicts, and kernel launch overhead.