parallel prefix scan

**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.