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.