Grid-Wide Barriers and Persistent Work Queues

A launch boundary is the only synchronisation primitive a kernel program has that is guaranteed to work at grid scope. Ending kernel A and starting kernel B orders every write from A before any read in B, on every device, with no cooperation from the program. The cost is a round trip through the host-side scheduler and a loss of everything that was resident: shared memory is uninitialised, registers are reallocated, and a grid that had warmed up its caches starts cold. A grid-wide barrier attempts to keep that ordering without the launch, and a persistent work queue builds on it to keep a single kernel resident for the whole application.

The barrier is a counter and a generation

The construction is forced by the hardware’s structure. Blocks do not share a scratchpad, so the counter that tracks arrivals has to live in global memory. Each block’s representative thread performs an atomic increment on the counter after its block’s own threads have arrived at a block barrier; then the block waits until the counter reaches the number of blocks.

Two details make it correct. The increment and the read must be ordered against the block’s other memory traffic. The release on the increment makes the block’s stores visible before the counter moves; the acquire on the observing read makes the stores of other blocks visible after. And the counter must be reset for the next barrier. The standard trick is a generation number: the counter’s value is compared against the expected total for the current generation, and a second barrier (or a sense-reversal per block) prevents a fast block from re-incrementing and being counted in the previous round.

__device__ unsigned arrive_count;
__device__ unsigned generation;

__device__ void grid_barrier(unsigned expected_blocks) {
    __syncthreads();                    // all threads of this block arrive
    if (threadIdx.x == 0) {
        // Release: this block's stores must be visible before it is counted.
        __threadfence();
        const unsigned gen = atomicAdd(&generation, 0u);   // relaxed read
        if (atomicAdd(&arrive_count, 1u) == expected_blocks - 1) {
            atomicExch(&arrive_count, 0u); // reset for the next round
            __threadfence();               // reset visible before the release
            atomicAdd(&generation, 1u);    // release the waiters
        } else {
            // The spin must re-read; a plain load may be hoisted out of it.
            while (atomicAdd(&generation, 0u) == gen) { }
            // Acquire: reads after this point see the other blocks' stores.
            __threadfence();
        }
    }
    __syncthreads();                    // all threads of this block resume
}

The fences are not decoration. Without the thread fence before the arrival increment, another block could observe the counter at its new value while this block’s data writes were still in flight, and read stale data through the barrier. The waiting side needs the mirror image: the load that observes the generation change must be an acquire, or the reads that follow it are not ordered after the observation. The same reasoning as the release-acquire pair described for counters applies, at grid scope rather than block scope, and it applies twice because there are two directions of handoff.

The residency requirement

The barrier deadlocks if a block that has arrived cannot make progress because the block it waits for has not been scheduled. Every block must be concurrently resident, or the barrier is only correct when the blocks that have arrived are exactly the ones still running — which is not a property the hardware guarantees.

The cooperative launch API exists to make that guarantee: the launch checks at dispatch time that the requested grid size fits the device’s resident capacity and fails rather than launching an overcommitted grid. That is the right failure, since a launch error is better than a hang, but it imposes a hard constraint on the grid size, which is why a persistent kernel does not size its grid to the problem. It sizes the grid to the machine, once, and the kernel loops over the problem in a grid-stride fashion. The stride loop described in grid-stride loops and launch size independence is a precondition for the persistent form, not an alternative to it.

Residency is a resource question, and it is easy to violate without noticing. Registers, shared memory per block, and the per-block resident limit all bear on how many blocks fit. A kernel compiled with higher register pressure than expected may reduce the achievable grid, and the cooperative check will report the maximum. On hardware with variable multiprocessor counts, the same binary may be launchable on one device and not another.

There is a second, less obvious requirement: the entire grid must be resident, not merely most of it. The maximum is a product of multiprocessor count, blocks per multiprocessor, and the resource limits that reduce the theoretical maximum, so a launch sized exactly at the computed limit is fragile: any reduction in capacity, whether from a concurrent kernel holding slots or from resources reserved by the driver, breaks it. A cooperative launch consequently excludes other kernels from running alongside it.

What the barrier buys

With a barrier available, a kernel can be structured as a sequence of phases with global state passed between them. Histogramming becomes one pass: build per-block histograms in shared memory, write them to global, barrier, then a subset of blocks combines them while the rest wait. Compaction becomes one pass: scan the predicate, barrier, scatter. Iterative solvers can check a convergence criterion in the middle of the kernel and stop early without a host round trip.

The last case is where the barrier earns its keep. A host-side loop that launches one kernel per iteration pays the launch latency and loses the resident state each time; an in-kernel loop with a barrier pays only the barrier. The convergence check becomes a reduction whose result is available to every block at the barrier, so the loop condition is a value in global memory read by all blocks after the barrier.

The cost is that all blocks wait for the slowest block at each barrier. A load imbalance that a launch boundary would absorb, because a block that finishes early simply retires and its slot is reused, becomes a stall for the whole grid. The longer the phases and the more variable their duration, the worse the barrier’s relative cost. That is the trade against splitting a workload into many tiny dependent launches: the barrier is cheaper than a launch, but it couples every block’s progress to the slowest one, and a workload whose per-phase durations vary widely can lose more to that coupling than it gains in dispatch time.

The work queue replaces the launch queue

A persistent kernel’s grid-stride loop distributes a fixed set of elements. A work queue distributes a variable set of tasks, which is what irregular workloads need: triangles of varying size, sparse matrix rows, or a recursive subdivision whose depth is data-dependent.

The simplest form is a device-side counter that each block’s leader atomically increments to claim the next task index, with the task list in global memory. Each block then processes the claimed task and comes back for another. The atomic is the bottleneck, because there is one contended location and one claim per task, so the claim is amortised by taking a range of indices rather than one. The block claims k consecutive tasks, processes them, and claims again. The cost per task falls by k and the tail at the end of the grid is at most k-1 wasted slots per block.

A more balanced form is the two-level queue: a small number of large tasks is handed out first, and each of those is subdivided dynamically into smaller tasks as blocks finish, so the grid drains evenly. The queue’s own state — head counter, tail counter, per-task completion flags — is exactly the machinery of the decoupled look-back scan, and the correctness conditions are the same: single-word state transitions, release-acquire ordering, and a chain that never waits on something that has not been scheduled.

Failure modes that only appear at this scale

Three defects are characteristic of persistent designs and rare in ordinary kernels.

A block that spins in a queue-pop loop without a forward-progress guarantee can wait for work that another block will enqueue only after its own queue pop. If the enqueuing block has not been scheduled — which cannot happen in a resident grid — the wait would be infinite; in a resident grid the wait terminates because every block is running. But the sequence of dependencies within the barrier structure must still be acyclic, and a queue whose consumers can also produce is where cycles appear.

The barrier’s reset logic is a race if the generation is incremented by a thread other than the one that reset the count. The single-writer form above is safe because the last arriving thread does both, with a fence between them and the release after, so no arrival can be counted twice or lost; a design that resets the counter in a separate phase or by a different block allows an arrival from the next round to be counted in the previous.

Memory ordering at grid scope is not implied by the barrier’s synchronisation property if the implementation uses relaxed atomics for the counter. The barrier is a synchronisation point only when the operations around it carry the appropriate ordering; a barrier built from relaxed atomics orders the blocks’ progress but not their memory traffic, and the resulting bug is a data race that appears only when the traffic crosses cache lines in an unlucky order.

When a launch is the better barrier

The launch boundary is free of these hazards and costs only the dispatch. For a pipeline of a handful of kernels over a large data set, the dispatch latency is small against the work in each kernel, and the separate kernels can each be tuned independently — a different block size for the scan phase and the scatter phase, for instance, which a persistent kernel cannot do without duplicating both in one function.

The persistent form wins when the per-phase work is small enough that dispatch dominates, when the state between phases is large enough that reloading it is expensive, or when the number of iterations is not known until runtime. It loses when the work is irregular enough that the barrier’s all-blocks-wait property dominates, and it loses badly when the grid cannot be made resident on the smallest target device. A library that ships one binary has to handle the case where the cooperative launch fails, which means a fallback path with a different structure — the cost that is easy to leave out of the design and impossible to leave out of the delivery.