GPU Scan (Prefix Sum) is the parallel algorithm that computes cumulative sums or other associative operations across array elements — where inclusive scan produces [a0, a0+a1, a0+a1+a2, ...] and exclusive scan produces [0, a0, a0+a1, ...], achieving 400-800 GB/s throughput (50-70% of peak bandwidth) through hierarchical implementation using warp primitives (__shfl_up_sync) for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and multi-pass algorithms for large arrays, making scan essential for algorithms like stream compaction (removing elements), radix sort (computing output positions), and sparse matrix operations where scan appears in 30-60% of advanced GPU algorithms and proper implementation using warp-level primitives and minimizing global memory accesses determines whether applications achieve 100 GB/s or 800 GB/s throughput.

Scan Fundamentals:

• Inclusive Scan: output[i] = input[0] + input[1] + ... + input[i]; includes current element; natural for cumulative sums
• Exclusive Scan: output[i] = input[0] + input[1] + ... + input[i-1]; excludes current element; useful for computing positions
• Associative Operations: sum, max, min, product, AND, OR, XOR; order matters for scan; enables parallelization
• Applications: stream compaction, radix sort, sparse matrix, work distribution; fundamental building block

Warp-Level Scan:

• Shuffle Up: use __shfl_up_sync() in loop; 5 iterations for 32 threads; no shared memory; 2-5× faster than shared memory
• Code Pattern: for (int offset = 1; offset < 32; offset *= 2) { int temp = __shfl_up_sync(0xffffffff, val, offset); if (lane >= offset) val += temp; }
• Performance: 500-1000 GB/s; 60-80% of peak bandwidth; 2-5× faster than shared memory; no synchronization overhead
• Use Cases: scan within warp; building block for block and grid scans; critical optimization

Block-Level Scan:

• Two-Stage: warp scan → inter-warp scan → add offsets; optimal at each level; 300-600 GB/s
• Pattern: each warp scans independently; scan warp sums; add warp offsets to elements; requires two syncs
• Shared Memory: store warp sums (32-64 elements); minimal usage; enables high occupancy
• Performance: 300-600 GB/s; 40-60% of peak bandwidth; 2-5× faster than pure shared memory

Large Array Scan:

• Multi-Pass: first pass scans blocks independently; second pass scans block sums; third pass adds offsets; 400-800 GB/s
• Three-Kernel: kernel 1 scans blocks; kernel 2 scans block sums; kernel 3 adds offsets; no global sync needed
• Single-Kernel: use cooperative groups grid.sync(); simpler but requires all blocks fit on GPU; 400-800 GB/s
• Performance: 400-800 GB/s for large arrays (>1M elements); 50-70% of peak bandwidth; memory-bound

Optimization Techniques:

• Warp Primitives: always use __shfl_up for intra-warp; 2-5× faster than shared memory; 500-1000 GB/s
• Bank Conflict Avoidance: pad shared memory arrays; prevents bank conflicts; 10-30% improvement
• Coalesced Access: ensure coalesced memory reads/writes; 128-byte aligned; achieves 100% bandwidth
• Multiple Elements Per Thread: each thread processes multiple elements; reduces overhead; 20-50% improvement

Inclusive vs Exclusive:

• Inclusive: simpler implementation; natural for cumulative sums; output[i] includes input[i]
• Exclusive: useful for positions; output[i] = sum of elements before i; radix sort, compaction
• Conversion: exclusive = shift(inclusive, 1) with 0 at start; inclusive = exclusive + input; trivial conversion
• Performance: same performance; choice based on application needs; exclusive more common in algorithms

Segmented Scan:

• Concept: scan multiple independent segments; each segment scanned separately; useful for batched operations
• Implementation: CUB DeviceSegmentedScan; segment flags or offsets specify boundaries; 300-600 GB/s
• Use Cases: per-row scan in matrix, per-group aggregation; graph algorithms; sparse matrix operations
• Performance: 300-600 GB/s; 40-60% of peak; depends on segment sizes; small segments have overhead

Thrust Scan:

• API: thrust::inclusive_scan(d_in.begin(), d_in.end(), d_out.begin()); or thrust::exclusive_scan()
• Performance: 400-800 GB/s; 60-80% of hand-tuned; 1 line vs 100-200 for custom implementation
• Customization: custom operators supported; thrust::plus(), thrust::maximum(), custom functors
• Use Cases: rapid development, general-purpose scan; acceptable 20-40% performance gap

CUB Scan:

• API: cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, N);
• Performance: 400-800 GB/s; 70-90% of hand-tuned; 10-30% faster than Thrust; lower-level API
• Features: inclusive, exclusive, custom operators; segmented scan; flexible
• Use Cases: performance-critical scans; fine-grained control; production systems

Stream Compaction:

• Pattern: scan predicate flags; use scan result as output positions; write elements to compacted array
• Code: flags = predicate(input); positions = exclusive_scan(flags); if (flags[i]) output[positions[i]] = input[i];
• Performance: 300-600 GB/s; 40-60% of peak; scan is bottleneck; 2-3× faster than CPU
• Use Cases: removing elements, filtering; sparse matrix, graph algorithms; data preprocessing

Radix Sort Integration:

• Histogram: count elements per bin; 300-600 GB/s
• Scan: exclusive scan of histogram; computes output positions; 400-800 GB/s
• Scatter: write elements to sorted positions; 200-300 GB/s
• Performance: scan is 20-40% of radix sort time; critical for overall performance

Work Distribution:

• Pattern: scan work counts; use scan result to assign work to threads; load balancing
• Code: work_counts = compute_work(input); positions = exclusive_scan(work_counts); assign work based on positions
• Performance: 300-600 GB/s for scan; enables balanced work distribution; 20-50% improvement in irregular algorithms
• Use Cases: irregular parallelism, dynamic work assignment; graph algorithms; sparse operations

Hierarchical Patterns:

• Three-Level: warp → block → grid; optimal at each level; 400-800 GB/s
• Warp Level: __shfl_up_sync(); 500-1000 GB/s; no shared memory; no sync
• Block Level: shared memory for inter-warp; 300-600 GB/s; minimal shared memory usage
• Grid Level: multi-pass or cooperative groups; 400-800 GB/s; depends on array size

Performance Profiling:

• Nsight Compute: shows memory bandwidth, warp efficiency, occupancy; identifies bottlenecks
• Metrics: achieved bandwidth / peak bandwidth; target 50-70% for scans; memory-bound
• Bottlenecks: uncoalesced access, bank conflicts, excessive synchronization; optimize patterns
• Tuning: adjust block size, elements per thread, padding; profile to find optimal

Common Pitfalls:

• Shared Memory Only: not using warp primitives; 2-5× slower; always use __shfl_up for intra-warp
• Bank Conflicts: unpadded shared memory; 2-10× slowdown; add padding to avoid conflicts
• Uncoalesced Access: stride access patterns; 10-100× slowdown; ensure coalesced reads/writes
• Too Many Syncs: excessive __syncthreads(); 10-30% overhead each; minimize synchronization

Best Practices:

• Warp Primitives: always use __shfl_up for intra-warp scan; 2-5× faster than shared memory
• Hierarchical: warp → block → grid; optimal at each level; 400-800 GB/s
• Pad Shared Memory: avoid bank conflicts; 1-2 elements padding; 10-30% improvement
• Coalesced Access: ensure coalesced memory reads/writes; 128-byte aligned; achieves 100% bandwidth
• Profile: measure actual bandwidth; compare with peak; optimize only if bottleneck

Performance Targets:

• Warp Scan: 500-1000 GB/s; 60-80% of peak; 2-5× faster than shared memory
• Block Scan: 300-600 GB/s; 40-60% of peak; optimal for 256-512 threads
• Grid Scan: 400-800 GB/s; 50-70% of peak; for large arrays (>1M elements)
• Overall: 400-800 GB/s for large arrays; 50-70% of peak bandwidth; memory-bound

Real-World Applications:

• Stream Compaction: removing invalid elements; 300-600 GB/s; 40-60% of peak; used in rendering, physics
• Radix Sort: computing output positions; 400-800 GB/s; 20-40% of sort time; critical for performance
• Sparse Matrix: CSR format construction; 300-600 GB/s; 30-50% of construction time
• Work Distribution: load balancing irregular work; 300-600 GB/s; 20-50% improvement in irregular algorithms

GPU Scan (Prefix Sum) represents the essential parallel primitive for position computation — by using hierarchical implementation with warp primitives for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and multi-pass algorithms for large arrays, developers achieve 400-800 GB/s throughput (50-70% of peak bandwidth) and enable algorithms like stream compaction, radix sort, and sparse matrix operations where scan is fundamental building block and proper implementation using warp-level primitives determines whether applications achieve 100 GB/s or 800 GB/s throughput.');