6 The Scan pattern: Hillis-Steele

This chapter introduces scan, or prefix sum, as the parallel counterpart to reduction: instead of collapsing an array into a single value, it preserves every intermediate running total. It explains the two common forms, inclusive and exclusive scan, using simple running-balance examples and showing how the exclusive version is just the inclusive result shifted right with the identity element inserted at the front. The text also emphasizes that scan is widely useful in real programs, from cumulative sums in numerical libraries to graphics, sorting, filtering, and other GPU algorithms.

The chapter then frames scan as a parallel algorithm design problem, where the goal is to reduce depth while keeping work reasonable. A sequential scan is easy but inherently chained by dependencies, since each output depends on the previous one. Hillis-Steele breaks that chain by computing partial sums in rounds: in round 1 each element combines with its left neighbor, in round 2 with the element two positions left, then four, and so on. Using double buffering avoids read-write collisions between threads, and a barrier at the end of each round ensures that every thread sees only the previous round’s results. The chapter builds the algorithm’s correctness around an invariant: after round r, each position contains the sum of a window of up to 2^r elements ending there.

Finally, the chapter presents a CUDA implementation of Hillis-Steele, organized into setup, shared-memory loading, the doubling-stride loop, guarded parallel additions, and copying the results back to global memory. It shows why boundary checks are needed at the left edge, why synchronization must happen after each round, and why the algorithm is correct for arrays of any size, not just powers of two. The tradeoff is made explicit: Hillis-Steele is fast in terms of depth, needing only logarithmic rounds, but it performs extra work—roughly n log n additions instead of the sequential n - 1. The chapter closes by warning about common mistakes and by positioning Hillis-Steele as the simple, intuitive scan algorithm that prepares the reader for a more work-efficient approach later.

The Hillis-Steele scan in 3 rounds, with double buffering. Each round doubles its stride (1, 2, 4) and writes results to the alternate buffer. Dotted boxes are cells left unchanged, dashed boxes are cells copied across, and solid boxes are newly computed sums. The pin (read source) and pout (write target) roles swap every round, so no thread overwrites a value another thread still needs. After round 3, buf[1] holds the final inclusive prefix sum.

Round 0 (the staging step). Each thread copies one element from the input array in global memory into its slot in buf[0] (shared memory). buf[1] is allocated but still uninitialized (?). No additions happen yet: this is just the starting state before the rounds begin.

Round 1 (stride = 1). buf[0] is the read source (pin) and buf[1] is the write target (pout). Each position i ≥ 1 reads buf[0][i] and buf[0][i−1], adds them, and writes the sum to buf[1][i]. Position 0 has no left neighbor at this stride, so its value is simply copied across (the leftmost dashed cell of buf[1]).

Round 2 (stride = 2). The roles flip: buf[1] becomes the read source (pin) and buf[0] becomes the write target (pout). Each position i ≥ 2 adds the value 2 positions to its left. Positions 0 and 1 are simply copied across: they already hold the correct prefix sums for their ranges.

Round 3 (stride = 4). The roles flip again: buf[0] is pin, buf[1] is pout. Each position i ≥ 4 adds the value 4 positions to its left. Only the 4 rightmost positions get updated; the first 4 were already correct from earlier rounds. After this final round, every position holds the correct inclusive prefix sum: [3, 4, 11, 11, 15, 16, 22, 25].

Case 1 (i ≥ 2ʳ − 1): full window.

Case 2 (2ʳ⁻¹ ≤ i < 2ʳ − 1): partial left window.

Case 3 (i < 2ʳ⁻¹): no left neighbor.

Round 1 (stride = 1) with cells classified by case. Position 0 is Case 3 (dotted): no left neighbor at stride 1, so its value is copied through. Positions 1 through 7 are Case 1 (solid): each has a full 2-element window. After this round, buf[1] holds sums of up to 2 consecutive elements at each position.

Round 2 (stride = 2) with cells classified by case. Positions 0 and 1 are Case 3 (dotted): they copy through unchanged. Position 2 is Case 2 (dashed): its left neighbor at position 0 holds only 1 element. Positions 3 through 7 are Case 1 (solid) with full 4-element windows.

Round 3 (stride = 4) with cells classified by case. Positions 0 through 3 are Case 3 (dotted), copying through. Positions 4 through 6 are Case 2 (dashed): their left neighbors hold partial sums reaching position 0. Only position 7 is Case 1 (solid) with a full 8-element window. The result, [3, 4, 11, 11, 15, 16, 22, 25], is the complete inclusive prefix sum.

• A scan computes every running total over an array, unlike a reduction, which returns only the final aggregate. Scan is the more general of the two: the last element of an inclusive scan is the reduction result, so a scan gives you the reduction for free, whereas a reduction alone cannot recover the running totals.
• Scan comes in two flavors: an inclusive scan includes the current element in its running total, while an exclusive scan excludes it. You can turn an inclusive result into an exclusive one by shifting it right by one position and inserting the operator’s identity element at the front.
• Total work counts the operations across all threads; depth counts the sequential rounds. The sequential scan does n−1 work in n−1 steps, and the goal of a parallel scan is to shrink the depth to O(log n) without letting the total work grow too far.
• The Hillis-Steele algorithm is a simple parallel scan that performs O(n log n) total work in O(log n) parallel time. It is easy to write and reason about, but it is not work-efficient, and whether the extra work matters for a given array size has to be measured.
• The algorithm is correct because of an invariant it preserves every round: after round r, each position holds the sum of the most recent (up to) 2ʳ input elements ending at itself. Once that window covers the whole array, every position reaches all the way back to position 0, which is the inclusive prefix sum.
• The kernel reads from one buffer and writes to the other, swapping the roles each round. This double buffering guarantees that no thread reads a value that has already been overwritten in the current round.
• A single __syncthreads() at the end of each round keeps all threads in step; the pin/pout swap is just two scalar assignments and needs no barrier of its own.
• Hillis-Steele trades extra work for speed. When the O(n log n) bill matters, the two-phase Blelloch algorithm in chapter 7 brings the total work down to O(n).
• Scan shows up throughout real code: NumPy’s np.cumsum(), PyTorch’s torch.cumsum(), radix sort, and stream compaction all rest on it; chapter 7 walks through them.

FAQ