🔗 Quiz: Parallel Patterns 25 XP Hard

In tree-based parallel reduction, each step combines pairs of values, halving the number of remaining elements: N → N/2 → N/4 → ... → 1. This takes log₂(N) steps. For N=1024, that's only 10 steps instead of 1023 serial additions. The work complexity is still O(N) (same total additions), but the step complexity drops from O(N) to O(log N), which is the entire point of parallelizing reduction.

With interleaved addressing (stride=1,2,4,...), the active threads access addresses that are spread apart — e.g., threads 0,2,4,6 access sdata[0],sdata[2],sdata[4],sdata[6] — causing 2-way bank conflicts. With sequential addressing (stride=N/2,N/4,...,1), active threads are contiguous (threads 0,1,2,...,stride-1) and access consecutive shared memory locations sdata[tid] and sdata[tid+stride]. The first access is conflict-free, and the second accesses a contiguous block offset by stride. This eliminates the bank conflicts that plague the interleaved version. Both approaches do the same number of additions.

In the standard tree reduction with sequential addressing, the stride starts at blockDim.x/2 and halves each round. At each step, threads with index < stride perform the addition: sdata[tid] += sdata[tid + stride]. After all log₂(256) = 8 rounds, only thread 0 remains active, and the final sum sits in sdata[0]. Thread 0 is typically responsible for writing this block-level partial sum to global memory. A second kernel or atomic operation then combines all block results.

__shfl_down_sync(mask, val, delta) is a warp-level intrinsic where each lane receives the value from the lane that is 'delta' positions higher (lane + delta). For example, lane 0 gets the value from lane delta, lane 1 from lane 1+delta, etc. Lanes near the end of the warp (where lane+delta ≥ warpSize) get their own value unchanged. This enables intra-warp reduction WITHOUT shared memory or __syncthreads() — just a sequence of __shfl_down_sync calls with delta = 16, 8, 4, 2, 1. It's faster because register-to-register communication is cheaper than shared memory round-trips.

Inclusive scan includes the current element in the running total: out[i] = data[0] + data[1] + ... + data[i]. So [3, 3+1, 3+1+4, 3+1+4+1, 3+1+4+1+5] = [3, 4, 8, 9, 14]. Exclusive scan excludes the current element: out[i] = data[0] + ... + data[i-1], with out[0] = 0 (the identity element). So [0, 3, 3+1, 3+1+4, 3+1+4+1] = [0, 3, 4, 8, 9]. Exclusive scan is critical for computing write offsets in stream compaction — out[i] tells you WHERE to write element i.

Blelloch's scan has two phases: an up-sweep (reduce) and a down-sweep, performing O(N) total additions across O(2 log N) steps. Hillis-Steele is simpler — each step, every element adds a value from distance 2^step behind — but ALL N elements do work at every step, totaling O(N log N) additions across O(log N) steps. Hillis-Steele is 'step-efficient' (fewer steps, same as Blelloch) but not 'work-efficient'. For GPUs where you have enough parallelism, Hillis-Steele's simplicity can sometimes win despite extra work, but Blelloch scales better for large arrays.

Blelloch's algorithm has two distinct phases: (1) Up-sweep: a standard parallel reduction that builds partial sums in a tree — after this phase, the last element contains the total sum. (2) Down-sweep: the last element is replaced with 0 (identity), then the tree is traversed top-down, distributing partial sums to produce the exclusive scan. Each phase takes O(log N) steps with O(N) total operations across both. This two-phase structure is why it's called a 'reduce-then-scan' pattern.

When thousands of threads from many blocks all call atomicAdd on the same global memory bins, the atomic unit serializes access — only one thread at a time can update a given bin. For skewed distributions where most values map to a few bins, this contention is catastrophic. The solution is privatized histograms: each block builds its own histogram in shared memory (only ~256 threads compete), then a single pass merges per-block histograms into global memory. This reduces contention by a factor of (total_threads / block_size).

Privatized histograms allocate one histogram copy per block in shared memory. Phase 1: each thread in a block uses atomicAdd on the shared-memory histogram — contention is limited to blockDim.x threads (e.g., 256) instead of the entire grid. Phase 2: threads in each block add their shared-memory histogram bins into the global histogram using atomicAdd. Since phase 2 has only (num_blocks × num_bins) atomics instead of (total_elements) atomics, contention drops dramatically. Option A (per-thread in registers) is possible for very small bin counts but impractical for typical histograms.

In tiled stencil computation, each block loads a tile of data into shared memory. But threads at the tile boundary need neighbor values that belong to adjacent tiles. Halo cells are the extra border elements loaded from neighboring tiles into shared memory. For a radius-1 stencil (5-point), each tile of size T×T must load (T+2)×(T+2) elements — the extra ring of width 1 around the tile is the halo. Without halo loading, boundary threads would need slow global memory reads for every neighbor access, defeating the purpose of tiling.

Stream compaction takes an input array and a predicate (e.g., x > 0), and produces an output array containing only elements that pass the predicate, packed contiguously with no gaps. The algorithm: (1) Evaluate the predicate → produce a flags array of 0s and 1s. (2) Exclusive prefix scan on flags → produces write indices. (3) Scatter: if flags[i]==1, write input[i] to output[scan[i]]. The exclusive scan is essential because scan[i] gives the exact output position for element i. This pattern is fundamental in ray tracing (compacting active rays), particle simulations, and sparse computations.

These are three fundamental data-parallel patterns. Map: out[i] = f(in[i]) — each element transformed independently, embarrassingly parallel. Gather: out[i] = in[index[i]] — reads from irregular (indexed) locations, writes contiguously. This is cache-unfriendly for reads but coalesced for writes. Scatter: out[index[i]] = in[i] — reads contiguously, writes to irregular locations. Scatter can have write conflicts if two elements map to the same index (solved with atomics). On GPUs, gather is generally preferred over scatter because coalesced writes are easier to achieve than coalesced reads.

For arrays larger than one block, a single block-level scan is insufficient. The reduce-then-scan (also called 'scan-then-propagate') approach: (1) Each block performs a local scan of its chunk and writes the block's total sum to an auxiliary array. (2) The auxiliary array of block totals is scanned (recursively if needed). (3) Each block adds the scanned block-total prefix to all its elements. This produces a correct global scan in O(N) work and O(log²N) steps. Libraries like CUB implement this pattern with optimizations like decoupled look-back to achieve single-pass performance.

Work complexity measures the total number of operations (analogous to sequential time). Step complexity measures the number of parallel rounds (depth of the computation DAG). Hillis-Steele: O(N log N) work, O(log N) steps — it's step-efficient (minimum parallel depth) but does extra redundant work. Blelloch: O(N) work, O(2 log N) steps — it's work-efficient (matches sequential cost) but has slightly higher depth constant. On a GPU with limited parallelism (fewer processors than N), Blelloch is usually better because extra work translates to real time. With unlimited processors, Hillis-Steele's lower step count wins.

CUB (CUDA UnBound) provides reusable, block-level and device-level primitives like cub::BlockReduce, cub::BlockScan, cub::DeviceReduce, and cub::DeviceHistogram. It gives developers fine-grained control over shared memory usage, tile sizes, and algorithm selection. Thrust is a higher-level, STL-like parallel algorithms library — you call thrust::reduce(d_vec.begin(), d_vec.end()) and it handles kernel configuration, memory management, and algorithm selection automatically. Both are shipped with the CUDA Toolkit. CUB is ideal when you need to embed scan/reduce inside your own kernels; Thrust is ideal for rapid prototyping and when performance tuning is less critical.