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.