gpu reduction patterns

**GPU Reduction Patterns** are **the parallel algorithms for combining array elements into single value through associative operations** — where hierarchical reduction using warp primitives (__shfl_down_sync) for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and atomic operations for inter-block (200-400 GB/s) achieves 60-80% of peak memory bandwidth and 2-10× speedup over naive implementations, making reduction optimization critical for applications like sum, max, min, dot product that appear in 40-80% of GPU kernels and proper implementation using warp-level primitives instead of shared memory, minimizing synchronization, and hierarchical patterns determines whether reductions achieve 100 GB/s or 1000 GB/s throughput. **Reduction Fundamentals:** - **Associative Operations**: sum, max, min, product, AND, OR, XOR; order doesn't affect result; enables parallelization - **Commutative**: most reductions also commutative; further optimization opportunities; non-commutative requires careful ordering - **Tree Pattern**: log2(N) steps to reduce N elements; each step halves active threads; optimal work complexity O(N) - **Memory Bound**: reductions are memory-bound; limited by bandwidth (1.5-3 TB/s); not compute; optimize memory access **Warp-Level Reduction:** - **Shuffle Down**: use __shfl_down_sync() in loop; 5 iterations for 32 threads; no shared memory; 2-10× faster than shared memory - **Code Pattern**: for (int offset = 16; offset > 0; offset /= 2) { val += __shfl_down_sync(0xffffffff, val, offset); } - **Performance**: 500-1000 GB/s; 60-80% of peak bandwidth; 2-5× faster than shared memory; no synchronization overhead - **Use Cases**: reduce within warp; building block for block and grid reductions; critical optimization **Block-Level Reduction:** - **Two-Stage**: warp reduction → shared memory → warp reduction; optimal at each level; 300-600 GB/s - **Pattern**: each warp reduces to single value; write to shared memory; first warp reduces shared memory; single thread has result - **Shared Memory**: 32-64 elements in shared memory (one per warp); minimal usage; enables high occupancy - **Performance**: 300-600 GB/s; 40-60% of peak bandwidth; 2-5× faster than pure shared memory **Grid-Level Reduction:** - **Three-Stage**: warp reduction → block reduction → atomic or multi-kernel; 200-400 GB/s - **Atomic Approach**: each block atomics final result; simple but contention; 200-400 GB/s for low contention - **Multi-Kernel**: first kernel reduces to per-block results; second kernel reduces blocks; no contention; 300-600 GB/s - **Cooperative Groups**: grid.sync() enables single-kernel multi-block; 200-400 GB/s; simpler than multi-kernel **Optimization Techniques:** - **Warp Primitives**: always use __shfl for intra-warp; 2-10× faster than shared memory; 500-1000 GB/s - **Minimize Sync**: reduce synchronization points; use warp primitives (no sync needed); 20-40% improvement - **Coalesced Access**: ensure coalesced memory reads; 128-byte aligned; achieves 100% bandwidth - **Multiple Elements Per Thread**: each thread processes multiple elements; reduces overhead; 20-50% improvement **Unrolling and Specialization:** - **Loop Unrolling**: unroll reduction loops; reduces overhead; 10-20% speedup; #pragma unroll - **Template Specialization**: specialize for block size; enables compile-time optimization; 10-30% speedup - **Warp Unrolling**: fully unroll warp reduction (5 iterations); eliminates loop overhead; 10-20% speedup - **Compile-Time Constants**: use template parameters for sizes; enables aggressive optimization; 20-40% improvement **Multiple Accumulators:** - **Pattern**: use multiple accumulators to increase ILP; reduces dependency chains; 30-60% speedup - **Code**: float sum1 = 0, sum2 = 0; for (int i = tid; i < N; i += stride) { sum1 += data[i]; sum2 += data[i+offset]; } - **Merge**: combine accumulators at end; single reduction; amortizes overhead - **Performance**: 30-60% improvement; exposes instruction-level parallelism; hides latency **Reduction with Transformation:** - **Pattern**: transform elements during reduction; fused operation; eliminates temporary storage - **Example**: sum of squares: reduce(x * x); dot product: reduce(a * b); L2 norm: sqrt(reduce(x * x)) - **Performance**: 2-5× faster than separate transform and reduce; eliminates memory traffic; 500-1000 GB/s - **Implementation**: thrust::transform_reduce() or custom kernel with inline transformation **Segmented Reduction:** - **Concept**: reduce multiple independent segments; each segment reduced separately; useful for batched operations - **Implementation**: CUB DeviceSegmentedReduce; segment offsets specify boundaries; 200-400 GB/s - **Use Cases**: per-row sum in matrix, per-group aggregation; graph algorithms; database operations - **Performance**: 200-400 GB/s; 40-60% of peak; depends on segment sizes; small segments have overhead **Thrust Reduce:** - **API**: float sum = thrust::reduce(d_vec.begin(), d_vec.end(), 0.0f, thrust::plus()); - **Performance**: 500-1000 GB/s; 70-90% of hand-tuned; 1 line vs 50-100 for custom implementation - **Customization**: custom operators supported; thrust::maximum(), custom functors - **Use Cases**: rapid development, general-purpose reduction; acceptable 10-30% performance gap **CUB Reduce:** - **API**: cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, N); - **Performance**: 500-1000 GB/s; 80-95% of hand-tuned; 10-20% faster than Thrust; lower-level API - **Features**: sum, min, max, custom operators; segmented reduce; flexible - **Use Cases**: performance-critical reductions; fine-grained control; production systems **Atomic Reduction:** - **Pattern**: each thread/warp/block atomics to global result; simple but contention-prone - **Warp Aggregation**: reduce within warp first; lane 0 atomics; 32× fewer atomics; 5-20× speedup - **Performance**: 200-400 GB/s with warp aggregation; 10-50 GB/s without; contention limits performance - **Use Cases**: simple reductions, low contention; when multi-kernel overhead unacceptable **Hierarchical Patterns:** - **Three-Level**: warp → block → grid; optimal at each level; 300-600 GB/s - **Warp Level**: __shfl_down_sync(); 500-1000 GB/s; no shared memory; no sync - **Block Level**: shared memory for inter-warp; 300-600 GB/s; minimal shared memory usage - **Grid Level**: atomic or multi-kernel; 200-400 GB/s; depends on contention **Performance Profiling:** - **Nsight Compute**: shows memory bandwidth, warp efficiency, occupancy; identifies bottlenecks - **Metrics**: achieved bandwidth / peak bandwidth; target 60-80% for reductions; memory-bound - **Bottlenecks**: uncoalesced access, excessive synchronization, low occupancy; optimize patterns - **Tuning**: adjust block size, elements per thread, unrolling; profile to find optimal **Common Pitfalls:** - **Shared Memory Only**: not using warp primitives; 2-10× slower; always use __shfl for intra-warp - **Excessive Sync**: too many __syncthreads(); 10-30% overhead each; minimize synchronization - **Uncoalesced Access**: stride access patterns; 10-100× slowdown; ensure coalesced reads - **Single Accumulator**: not using multiple accumulators; limits ILP; 30-60% slower **Best Practices:** - **Warp Primitives**: always use __shfl for intra-warp reduction; 2-10× faster than shared memory - **Hierarchical**: warp → block → grid; optimal at each level; 300-600 GB/s - **Multiple Accumulators**: use 2-4 accumulators; increases ILP; 30-60% improvement - **Coalesced Access**: ensure coalesced memory reads; 128-byte aligned; achieves 100% bandwidth - **Profile**: measure actual bandwidth; compare with peak; optimize only if bottleneck **Performance Targets:** - **Warp Reduction**: 500-1000 GB/s; 60-80% of peak; 2-5× faster than shared memory - **Block Reduction**: 300-600 GB/s; 40-60% of peak; optimal for 256-512 threads - **Grid Reduction**: 200-400 GB/s; 30-50% of peak; limited by atomic contention or multi-kernel overhead - **Overall**: 500-1000 GB/s for large arrays (>1M elements); 60-80% of peak bandwidth **Real-World Applications:** - **Sum**: array sum, vector norm; 500-1000 GB/s; 60-80% of peak; critical building block - **Dot Product**: vector dot product; 500-1000 GB/s; fused multiply-add and reduce; 60-80% of peak - **Max/Min**: find maximum/minimum; 500-1000 GB/s; 60-80% of peak; used in normalization - **Statistics**: mean, variance, standard deviation; 400-800 GB/s; multiple reductions; 50-70% of peak GPU Reduction Patterns represent **the fundamental parallel primitive** — by using hierarchical reduction with warp primitives for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and atomic operations for inter-block (200-400 GB/s), developers achieve 60-80% of peak memory bandwidth and 2-10× speedup over naive implementations, making reduction optimization critical for GPU applications where reductions appear in 40-80% of kernels and proper implementation using warp-level primitives instead of shared memory determines whether reductions achieve 100 GB/s or 1000 GB/s throughput.