Parallel Prefix Scan Inclusive Exclusive

Parallel Prefix Scan is a foundational parallel algorithm computing cumulative sums (or generic associative operation) across array elements with logarithmic depth, essential for stream compaction, sorting, and GPU application performance.

Inclusive vs Exclusive Scan Definitions

• Inclusive Scan: Output[i] = sum(Input[0..i]). Last element = total sum. Example: input=[1,2,3,4] → output=[1,3,6,10].
• Exclusive Scan: Output[i] = sum(Input[0..i-1]). Last element = sum of first N-1 elements. Example: input=[1,2,3,4] → output=[0,1,3,6].
• Scan Generalization: Replace addition with generic associative operator (min, max, bitwise AND/OR). Inclusive/exclusive semantics apply to any operator.
• Prefix Operation: Generic term for all cumulative operations. Scan = GPU terminology; prefix sum = generic algorithmic terminology.

Blelloch Up-Sweep/Down-Sweep Algorithm

• Work-Efficient Scan: O(N) work (same as sequential), O(log N) depth (parallel). Contrasts with Hillis-Steele (O(N log N) work).
• Up-Sweep Phase: Tree reduction bottom-up. Pair adjacent elements (stride 1), add to right neighbor. Stride doubles (2, 4, 8, ...); continue log(N) levels.
• Down-Sweep Phase: Tree expansion top-down. Implicit left child, propagate partial sums. Builds final scan output.
• Example: Array [1,2,3,4,5,6,7,8]. Up-sweep computes sums at each level. Down-sweep produces [0,1,3,6,10,15,21,28] (exclusive scan).

Hillis-Steele Step-Efficient Scan

• Step Efficient: O(N log N) work, O(log N) depth. Work redundancy acceptable for parallelism benefit.
• Algorithm: Pass k adds elements at offset 2^(k-1). Pass 0 adds offset 1; pass 1 adds offset 2; pass 2 adds offset 4.
• Correctness: After log(N) passes, all elements have received all necessary contributions. Output = exclusive scan.
• Implementation: Simple, shorter code. Cache-efficient (stride pattern locality). Common GPU implementation (thrust::inclusive_scan).

GPU Scan Implementation Using Shared Memory

• Block-Level Scan: Shared memory scan for block_size elements (typically 512-1024). Load thread-local elements into shared memory, perform in-memory scan.
• Warp-Level Scan: Fast scan within warp (32 threads) using __shfl_sync. Shuffle broadcasts partial sums to other threads.
• Multi-Block Scan: Larger arrays require multiple blocks. Scan each block independently, compute block sums, broadcast block sums.
• Bank Conflict Avoidance: Access shared memory with stride patterns avoiding bank conflicts. Padding needed for certain strides.

Warp-Level Scan via Shuffle Operations

• __shfl_sync(): Warp communication intrinsic. Broadcasts register value from thread srcLane to all threads in warp.
• Shuffle-Based Scan: Warp 0 computes scan in place via shuffle operations (no shared memory). Each thread maintains partial sum, shuffles intermediate results.
• Efficiency: No shared memory contention. Low latency (3-5 cycles per shuffle). Throughput sufficient for 32 threads.
• Recursive Calls: Large arrays call shuffle-scan on successive 32-element segments, merge results via global memory.

Segmented Scan and Applications

• Segmented Scan: Multiple independent scans on contiguous segments. Segment boundaries determined by flag array.
• Implementation: Carry-in value propagates across segment boundaries. Boundary detected via flag; carry-in reset for new segment.
• Stream Compaction Use Case: Predicate array (element 'selected' or not) scanned. Output indices for selected elements computed via segmented exclusive scan.
• RLE Compression: Run-length encoding uses segmented scan to compute output positions for each run.

Prefix Sum Applications in GPU Computing

• Stream Compaction: Filter array removing unwanted elements. Boolean predicate scan yields output indices; compacted array built from selected elements.
• Radix Sort: Counting sort histogram (count elements in each bucket), then prefix scan of counts yields output positions. Elements scattered to output based on positions.
• Histogram Computation: Count occurrences of each value. Atomic-based histogram slow (contention). Segmented scan groups histogram updates.
• Polynomial Evaluation: Horner's method (y = a_n + x(a_{n-1} + x(a_{n-2} + ...))). Scan propagates intermediate values. Fine-grain parallelism via scan.

Performance and Scalability

• Blelloch Complexity: O(N) work, O(log N) depth. Practical speedup for 1000s of elements (overhead amortizes). For 100+ million elements, multiple GPU kernels necessary.
• Bandwidth: Scan memory-bound (two passes over data per scan). Peak bandwidth utilization ~50-80% typical. Bottleneck: memory, not computation.
• Throughput Library: Libraries (thrust::inclusive_scan, CUB) heavily optimized. Typical: 200-500 GB/s actual bandwidth (vs 2 TB/s peak HBM). Algorithm variants chosen per GPU model.
• Multi-GPU Scaling: Global scan across GPUs requires local scans → allreduce (all block sums) → broadcast (global offset) → add offset to local results. Allreduce dominates communication cost.