Single-Pass Prefix Scan with Decoupled Look-Back

An exclusive prefix sum turns an array into the array of sums of preceding elements: input [3, 1, 4, 1] becomes [0, 3, 4, 8] exclusive, or [3, 4, 8, 9] inclusive. It is the primitive behind stream compaction, radix sort digit offsets, run-length decoding and any transform that needs each element to know how much work precedes it. The parallel difficulty is that element $i$ depends on everything before it. The textbook answers are the Hillis-Steele form, which does $\log n$ passes of $n$ additions, and the Blelloch form, which reduces to a single value and then sweeps back up, doing $O(n)$ total additions over $2\log n$ levels. Neither maps cleanly onto a grid of blocks: the first costs more work than a sequential scan, and the second is efficient only at sizes that fit inside one block. The practical form is a two-level decomposition, and the interesting engineering is in how its two levels are joined.

The two-level decomposition and its extra pass

The straightforward approach: each block scans its own chunk in shared memory, writes the inclusive total to a per-block array, and a second kernel scans that array, and a third kernel adds each block’s exclusive offset back to its elements. Three launches, and the intermediate arrays are small — one entry per block — but the middle pass is a launch with a kernel that does almost no work per thread, and the whole sequence is serialised by launch boundaries.

The alternative folds the middle pass into the end of the first kernel with a grid-wide barrier, requiring every block to remain resident, which in turn requires the launch to be no larger than the machine can hold. That is the persistent-kernel constraint: a grid-stride loop over all the chunks, one barrier, then the same loop again. It works, and it costs occupancy freedom because the launch size can no longer exceed the device’s capacity.

Look-back removes both constraints. Blocks are launched normally, at whatever size, and each block waits only for the specific blocks whose totals it needs.

Three states in a per-block descriptor

Each block owns one slot in a global array of descriptors. The slot holds two values: an aggregate — the sum of the block’s entire chunk — and a prefix — the exclusive scan of everything before the block. The slot also has a flag, encoded in one of the values, with three states:

State	Meaning	What a reader may use
invalid	block has published nothing	nothing
aggregate available	block has its own total	the total, as a partial sum
prefix available	block has its exclusive prefix	the prefix, which is final

The block computes its local scan first, entirely in registers and shared memory, giving it its aggregate. It then publishes the aggregate with the “aggregate available” flag. Then it looks back: it reads the predecessor’s slot and accumulates backwards until it either reaches a slot with a prefix available or accumulates enough aggregates to cover the distance to block 0. If it reaches a prefix, the block’s own prefix is that value plus the accumulated partials, and it republishes its slot as a prefix. If it reaches block 0, the accumulated partials are the prefix.

The key property is that a block never has to wait for immediate predecessors to publish a prefix — only an aggregate. Aggregates are available almost immediately, because they are computed from data the block already has. A block that is far behind can consume a chain of aggregates from blocks that will publish their prefixes later, and the condition for progress is weak: as long as some block eventually publishes a prefix, everyone behind it can finish.

The look-back is a polling loop, so it burns issue slots and memory bandwidth while waiting. It terminates because the block with index 0 publishes a prefix as soon as it has its aggregate, which starts a cascade.

Why the work stays linear

A naive look-back reads every predecessor’s slot, so block $k$ performs $k$ reads and the total is $O(\text{blocks}^2)$. The decoupling is what avoids that: blocks publish aggregates as soon as they have them, so a block that looks back usually finds a prefix published by one of its predecessors and stops after one or a few reads. In the common case the cascade is nearly synchronous across blocks and each block performs a small constant number of reads.

The bound depends on scheduling. If the grid is small enough that all blocks are resident, aggregates appear in roughly launch order and the cascade propagates quickly. If the grid is much larger than the machine, later blocks are dispatched after earlier ones have finished, so when a later block starts, its predecessors have already published prefixes — the look-back is a single read. The bad case is a grid that is slightly larger than the resident set, where a wave of blocks starts together and each has to wait for the cascade to reach it.

There is a version of the algorithm with a bounded look-back window: instead of following the chain arbitrarily far, each block accumulates at most $W$ predecessors and, if it has not reached a prefix by then, publishes an intermediate state that records the accumulated sum and the distance. This caps the per-block work at $W$ reads and keeps the total linear with a larger constant. The window size is the tuning parameter: a wide window makes the common case one read but costs more reads in the worst case; a narrow window bounds the worst case and makes cascade propagation slower.

The publication protocol in code

The descriptor is a single 64-bit word holding a flag in the high bits and a value in the low bits, which allows the whole state transition to be a single atomic store. Writing the flag and the value as separate stores would allow a reader to observe the new value with the old flag, or vice versa.

// Descriptor: the state occupies the high 32 bits, the payload the low 32.
#define STATE_INVALID  0ull
#define STATE_AGGREGATE 1ull
#define STATE_PREFIX   2ull

__device__ unsigned long long* block_desc;  // one per block

__device__ void publish(int b, int state, unsigned sum) {
    // One store publishes the state and the value together.
    atomicExch(&block_desc[b], ((unsigned long long)state << 32) | sum);
}

The reader must not combine a fresh flag with a stale payload, so it reads the whole word atomically and parses it. A read that observes the prefix state has a final value for that block; a read that observes the aggregate state has a partial sum that can be accumulated but not relied on as final. Because the writer performs a single atomic exchange and the reader a single atomic load, the pair needs no additional ordering for the value itself; the ordering that matters is between the block’s shared-memory scan and its publication, which a block-level barrier handles.

The publication is also where the memory scope matters. A descriptor that only threads on the same device will read needs device scope; declaring system scope is correct but can cost more, because the operation must be ordered against traffic from the host and from peer devices that will never touch it. The payload must be visible to any thread that observes the flag, which in a single-word descriptor is guaranteed by the payload and flag being one atomic access rather than by any fence. Code that splits them into a value store and a separate flag store needs the release-acquire pair described in atomics and contention in counters and aggregates.

Where the algorithm meets the hardware

Three practical constraints shape an implementation.

The grid must be limited, because the descriptor array is one slot per block and the cascade’s latency grows with the number of blocks. A large grid means a long chain of look-backs in the worst case, and a long chain of dependent reads is exactly what a GPU handles poorly. Production implementations cap the grid and use a grid-stride loop so each block covers many chunks; the scan then has two levels again, but both are inside one launch.

The look-back is a spin, and a spin on modern hardware must use a load that does not hoist out of the loop. A compiler is entitled to assume that a location not written by the current thread does not change during a loop and may load it once, which turns the spin into an infinite loop. The load must be volatile or an atomic, and the refresh must be explicit on each iteration.

Forward progress is not guaranteed by the programming model in the way a lock-free algorithm on a CPU expects. A spinning block is not guaranteed to yield its slot, and if the block it waits for has never been scheduled, the wait never ends. The algorithm avoids this only because the cascade reaches its start: block 0 never waits, so progress propagates from the beginning. Any design in which a block waits for a block that might be dispatched later is deadlock-prone, and this is the reason look-back chains only backwards.

When not to use it

The look-back scan is the right choice when the scan is on a hot path with a large input and the extra launches of the two-level form are a measurable fraction of the time, or when the application cannot afford the extra global-memory traffic of the intermediate arrays. For a scan over a few hundred thousand elements the two-level form is often the faster choice, because it has no spinning and no descriptors, and the extra launch costs on the order of microseconds.

There is also a correctness condition that is easy to miss: the scan operator must be associative, and the algorithm’s accumulation order determines which grouping the results reflect. For integer addition that is invisible. For floating-point addition it means the result depends on how the look-back happened to group the aggregates, so the scan is not reproducible across runs — the same caveat as the tree reduction described in floating-point summation order, amplified by the non-determinism of which blocks publish first.

For compaction — the most common use — the scan is over a predicate count, which is integer, and the non-determinism does not exist. That is why the pattern is so closely associated with stream compaction and radix sort, where the values are counts and the ordering guarantee is exact.