parallel reduction algorithms
**Parallel Reduction Algorithms** — Techniques for combining a collection of values into a single result using an associative operator, executed across multiple processing elements simultaneously.
**Tree-Based Reduction Patterns** — Binary tree reduction pairs adjacent elements in each step, halving the active processors at every level to complete in O(log N) steps. The upward sweep phase combines partial results from leaves to root, while the downward sweep can distribute the final result back. Recursive doubling has each processor communicate with a partner at exponentially increasing distances, keeping all processors active but requiring more communication bandwidth. Butterfly reduction combines elements of both approaches, achieving optimal latency and bandwidth utilization for power-of-two processor counts.
**Work-Efficient Parallel Scan** — Blelloch's work-efficient scan algorithm performs a reduction in the up-sweep phase followed by a down-sweep that computes all prefix sums. The total work is O(N) matching the sequential algorithm, while the span is O(log N). Inclusive scan includes the current element in each prefix result, while exclusive scan shifts results by one position. Segmented scan extends prefix operations to operate independently within segments defined by flag arrays, enabling nested parallelism patterns.
**GPU-Specific Reduction Techniques** — Warp-level reductions use __shfl_down_sync() to exchange values between threads within a warp without shared memory, completing a 32-element reduction in 5 steps. Block-level reductions combine warp-level results through shared memory, with the first warp performing a final warp reduction. Grid-level reductions use atomic operations or multi-pass kernels where each block produces a partial result that subsequent kernels combine. Cooperative groups in CUDA enable flexible reduction scopes beyond fixed warp and block boundaries.
**Optimization Strategies** — Sequential addressing in shared memory reductions avoids bank conflicts that plague interleaved addressing patterns. Unrolling the last warp of a reduction eliminates unnecessary synchronization barriers since warp execution is inherently synchronous. Processing multiple elements per thread during the initial load phase reduces the number of active threads needed and improves arithmetic intensity. For non-commutative operators, maintaining element order requires careful indexing that preserves the original sequence during the tree traversal.
**Parallel reduction algorithms are among the most fundamental building blocks in parallel computing, serving as the basis for aggregation, prefix sums, and countless higher-level parallel patterns.**