Overview
5 The Reduction pattern
Reduction is the parallel pattern for turning many values into one result, such as a sum, maximum, minimum, product, or count. The central requirement is that the combining operation be associative, so the values can be regrouped without changing the answer; an identity element is also needed so extra threads can safely contribute “nothing” when there are more threads than data. This idea shows up everywhere, from spreadsheet-style totals and database aggregates to GPU tensor operations and large-scale frameworks for distributed data processing.
The chapter first shows why a straightforward parallel port fails. A single running total creates a loop-carried dependency in the sequential version, and making every thread update one shared variable leads to race conditions because read-modify-write is not atomic. Although atomic operations fix correctness, using one atomic update per thread serializes the work and largely destroys the benefit of parallelism. The better answer is to reduce in stages: each block computes a local subtotal, and then those partial results are merged safely.
The efficient solution is a tree-based reduction implemented inside a block using shared memory and synchronization barriers. Threads first load values from global memory into shared memory, padding out-of-range positions with the operation’s identity, then repeatedly combine pairs while the stride shrinks by half each round until one value remains in sdata[0]. Finally, one thread per block uses an atomic operation to add the block’s subtotal to the global result. The chapter closes by showing how the same structure adapts to max and min, warning about floating-point rounding differences, and highlighting common mistakes such as missing barriers, forgetting initialization, using non-atomic updates, or mismatching shared-memory size and block size.
A knockout tournament halves the number of teams every round until one champion remains; tree reduction does the same with numbers, halving the active elements each round (stride 4 → 2 → 1) until a single sum is left. Eight numbers collapse to one in just three rounds, instead of the seven steps a sequential sum would take.
The GPU memory hierarchy. Registers are the fastest; each thread has its own dedicated registers. All the threads in one block can read and write to the same shared memory. Global memory is reachable by every thread across all blocks, but it lives off-chip, which makes it by far the slowest.
The tree reduction kernel in three acts: load, reduce, write back. Act 1 copies global memory into shared memory so the fast reduction can happen on-chip. Act 2 runs three rounds of stride-based addition (stride 4, 2, then 1), collapsing eight values down to one. Act 3 uses atomicAdd to safely write the block's result (36) back to global memory.
Act 1 (Load): staging data in shared memory for fast access. Each thread reads one value from slow global memory into its private register, then writes it to shared memory, a two-step relay that happens in parallel across all eight threads simultaneously. Once sdata is populated, the reduction can run entirely on-chip without touching global memory again.
Act 2 (Reduce): three rounds of pairwise addition in shared memory. Each round, active threads read two values from sdata into registers, add them, and write the result back, halving the number of active threads each time (4 → 2 → 1). By the end of Round 3, the block's total sum (36) sits in sdata[0], ready to be written back to global memory.
Act 3 (Write back): one thread per block delivers the result. Thread 0 reads the block's sum from sdata[0] into a register, while T1–T7 sit idle. It then uses atomicAdd to safely add that value to the global result, necessary because multiple blocks may be writing to the same location simultaneously.
atomicAdd guarantees the right answer, regardless of which block goes first. Block 0 and Block 1 both race to update the same result in global memory with their partial sums (36 and 24). Whether Block 0 arrives first (0 → 36 → 60) or Block 1 does (0 → 24 → 60), the final sum lands on 60; atomicAdd serializes the updates so no contribution is ever lost.
Why atomicAdd is safe: read-modify-write as one indivisible step. When two blocks target the same address, the hardware serializes their requests: while one block's three steps (read → modify → write) run, the others wait, and the next one begins by reading the updated value, so no update is lost. When blocks target different addresses, their read-modify-writes run fully in parallel. The order in which same-address requests are served is unpredictable, but the final sum is always correct.
End-to-end reduction of 16M floats on a Tesla T4 (transfers included). CPU 52 ms; GPU-atomic 47.5 ms (1.10× over CPU); GPU-tree 14.8 ms (3.52×). The basis is end-to-end on purpose: the atomic bar sitting almost level with the CPU bar is the point. A direct GPU port of the sequential loop barely beats the CPU. Measured on a Google Colab Tesla T4 with CUDA 12.8 (nvcc -O3), 2²⁴ single-precision values, 256 threads per block; the text explains what each timing covers. The numbers compare the three approaches on one machine, not the speedup to expect elsewhere.
Summary
- The Reduction pattern combines many values into one using an associative operation.
- Naive parallel reduction creates race conditions and gives unpredictable, wrong results.
- Tree-based reduction takes O(log n) depth (the number of rounds that must run one after another) instead of O(n), enabling large speedups when enough threads are available.
- Shared memory is a fast, per-block scratch space for thread communication.
- Padding out-of-bounds threads with the identity element keeps partial blocks correct, and the halving stride needs a power-of-2 block size.
- __syncthreads() is a barrier that synchronizes all threads in a block; skipping one is the most common reduction bug.
- Atomic operations prevent race conditions when combining results across blocks.
- Structure matters more than hardware: on 16M floats the atomic-per-element version only ties the CPU, while the tree reduction runs 3.52× faster.
- The same tree code works for other associative operations: swap the operator and its identity element to get max, min, product, or count, provided the type and the final atomic (or a second pass) support that operation.
- NumPy’s sum(), PyTorch’s max(), and SQL’s COUNT()are all examples of the Reduction pattern at work.
- The same idea shows up across the parallel family: MPI_Allreduce combines partial sums across a cluster, and OpenMP’s reduction clause does it on a multicore CPU.
Grokking Parallel Programming ebook for free
