Parallel Prefix Sum (Scan) is the fundamental parallel computing primitive that computes all prefix sums of an input array — where prefix_sum[i] = sum(input[0..i]) for all i simultaneously — in O(N/P + log P) time on P processors, serving as the building block for parallel sorting, stream compaction, histogram computation, and virtually every parallel algorithm that requires computing cumulative quantities or assigning output positions dynamically.

Why Scan Is Foundational

Sequentially, computing prefix sums is trivial: iterate once, accumulating. But in parallel, each element's result depends on all preceding elements — a seemingly inherently serial dependency chain. The breakthrough is that associative operations can be restructured into a tree pattern that computes all prefixes in O(log N) parallel steps.

Scan Variants

• Exclusive Scan: output[i] = sum(input[0..i-1]). output[0] = 0 (identity element). Used for computing output positions (scatter addresses).
• Inclusive Scan: output[i] = sum(input[0..i]). Each element includes itself.
• Segmented Scan: Scan that resets at segment boundaries. Enables parallel processing of variable-length sequences (e.g., per-row operations in sparse matrices).

Blelloch's Work-Efficient Parallel Scan

Two phases on a balanced binary tree: 1. Up-Sweep (Reduce): Bottom-up, each node stores the sum of its left and right children. After log₂(N) steps, the root contains the total sum. Total work: N-1 additions. 2. Down-Sweep: Top-down, the root receives 0 (identity). Each node passes its value to its left child and (its value + left child's old value) to its right child. After log₂(N) steps, every leaf contains its exclusive prefix sum. Total work: N-1 additions.

Total work: 2(N-1) additions — the same as sequential (work-efficient). Span: 2×log₂(N) steps.

GPU Implementation

• Block-Level Scan: Each thread block loads a tile (512-2048 elements) into shared memory and performs the up-sweep/down-sweep within the block using __syncthreads() barriers.
• Grid-Level Scan: For arrays larger than one block's capacity, a three-kernel approach: (1) block-level scan producing per-block partial sums, (2) scan of the partial sums, (3) add each block's prefix to all elements in that block.
• Warp-Level Scan: For the last 32 elements, use __shfl_up_sync() to perform prefix sum within a warp without shared memory or barriers — the fastest possible implementation.

Critical Applications

• Stream Compaction: Given an array and a predicate, extract elements that satisfy the predicate into a dense output array. Scan of the predicate flags computes the output position for each surviving element.
• Radix Sort: Each pass of parallel radix sort uses scan to compute the destination index for each element based on its digit value.
• Sparse Matrix Construction: Scan of row nonzero counts computes the row pointer array (CSR format) for a sparse matrix.
• Histogram: Scan of bin counts computes cumulative histogram (used in histogram equalization and percentile computation).

Parallel Scan is the universal addressing primitive of parallel computing — the algorithm that answers "where does each element go?" in parallel, enabling the dynamic data movement that makes complex parallel algorithms possible on architectures with no shared mutable state.