parallel prefix scan
**Parallel Prefix Sum (Scan)** is the **fundamental parallel primitive that computes all prefix sums of an input array — where output[i] = input[0] + input[1] + ... + input[i] for inclusive scan — in O(N/P + log P) time on P processors, serving as the building block for stream compaction, radix sort, histogram, sparse matrix operations, and dozens of other parallel algorithms that require computing cumulative results across data**.
**Why Scan Is a Fundamental Primitive**
Sequentially, prefix sum is trivial: a single loop accumulating values. But in parallel, every output depends on all previous inputs — an apparently serial dependency. The breakthrough insight (Blelloch, 1990) is that this dependency can be resolved in O(log N) parallel steps using a two-phase (up-sweep/down-sweep) algorithm, making scan the most important parallel building block after reduction.
**Blelloch's Work-Efficient Scan**
**Phase 1: Up-Sweep (Reduce)**
```
Input: [3, 1, 7, 0, 4, 1, 6, 3]
Step 1: [3, 4, 7, 7, 4, 5, 6, 9] (pairs summed)
Step 2: [3, 4, 7, 11, 4, 5, 6, 14] (stride-4 sums)
Step 3: [3, 4, 7, 11, 4, 5, 6, 25] (total sum at last position)
```
**Phase 2: Down-Sweep (Distribute)**
```
Set last to 0, then distribute partial sums back down the tree.
Result: [0, 3, 4, 11, 11, 15, 16, 22] (exclusive prefix sum)
```
Total work: O(2N) — same as sequential. Depth: O(2 log N). Work-efficient, unlike the naive algorithm that does O(N log N) work.
**GPU Implementation**
1. **Block-Level Scan**: Each thread block loads a tile (e.g., 1024 elements) into shared memory and performs the up-sweep/down-sweep within the block using __syncthreads() barriers.
2. **Block Sum Extraction**: Each block's total sum is stored in an auxiliary array.
3. **Block Sum Scan**: A recursive scan computes prefix sums of the block totals.
4. **Final Propagation**: Each block adds its block-level prefix to all its elements, producing the globally correct prefix sum.
NVIDIA CUB and Thrust provide highly-optimized scan implementations achieving >90% of peak memory bandwidth.
**Applications**
- **Stream Compaction**: Given an array and a predicate, produce a new array containing only elements satisfying the predicate. Scan computes the output indices: scan the predicate array, each true element writes to the index given by its scan value.
- **Radix Sort**: Each radix pass uses scan to compute scatter positions for each digit bucket.
- **Sparse Matrix**: CSR format uses scan over row lengths to compute row pointer offsets.
- **Histogram**: Scan over bin counts produces cumulative distribution functions.
- **Dynamic Work Generation**: Scan determines output offsets when each input produces a variable number of outputs (e.g., each triangle produces 0-N fragments in rasterization).
**Parallel Prefix Scan is the secret weapon of parallel algorithm design** — the primitive that converts apparently sequential cumulative computations into fully parallel operations, enabling efficient GPU implementations of algorithms that would otherwise resist parallelization.