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.