Block Barriers and Divergent Synchronization

A block is the only scope in which a kernel can both communicate through shared memory and synchronise without leaving the kernel. The barrier that makes this possible is one instruction with two jobs: it waits for every thread in the block to arrive, and it orders the memory operations issued before it against the ones issued after it. Both jobs are done inside a single multiprocessor, which is why the barrier costs tens of cycles in the common case rather than the thousands a cross-device synchronization would cost.

Two guarantees in one instruction

__syncthreads() behaves like a combined barrier and memory fence, restricted to the block. In PTX it lowers to a bar.sync instruction, which suspends the executing thread until the stated number of threads in the block have arrived at that barrier, then releases them together. The count is an operand, which is why the named-barrier forms can synchronise a subset. The memory ordering part is what allows the classic pattern: thread $t$ writes tile[t], the barrier, then thread $t$ reads tile[blockDim.x - 1 - t]. Without the ordering, the reads might be issued before the writes have landed, because the compiler and the hardware both reorder memory operations to the same storage.

The scope is the block, and this is a hard partition of the machine’s memory behaviour. Two threads in the same block share an L1 and a shared-memory partition; a barrier can therefore be a counter operation with no traffic leaving the multiprocessor. Threads in different blocks may be on different multiprocessors and there is no counter that both can see, which is why there is no __syncthreads() that spans a grid. The mechanisms that do span a grid are described in grid-wide barriers and persistent work queues.

The ordering guarantee covers shared memory and global memory. In practice the interesting case is shared memory, because a global-memory round trip through the same barrier is unusual — a block has no reason to write to global memory and read its own writes back except in the double-buffer patterns where both copies live in global memory.

Why partial arrivals are undefined

The count-based design is the source of the restriction. Suppose a block of 256 threads contains a branch that puts a barrier on each side:

if (threadIdx.x < 128) {
    __syncthreads();
} else {
    __syncthreads();
}

Each side issues one barrier, so the source contains one arrival per thread and the program looks as though it makes progress. It does not. A warp that straddles the condition runs one side first, with only the lanes that took that branch active, and those lanes wait at the barrier for the block’s full count. The lanes on the other side of the branch are parked behind it and cannot run until the first side completes, so the arrivals they would contribute never reach this barrier instance at all. The block hangs until the launch is killed.

The failure is not always a hang. Divergence is dynamic, so a warp whose lanes all agree on the condition executes one side and skips the other entirely, and a warp that disagrees executes both in sequence. Which warps fall into which category depends on the thread index mapping, and a block whose warps happen to be uniform reaches the barrier once with all 256 arrivals and proceeds normally. The same source, with a different block size, deadlocks. That sensitivity to a launch parameter is why the language defines the construct as undefined rather than as a hang, and why code that happens to work is not evidence that it is correct.

There is a corollary that surprises people who come from CPU threading: a barrier inside a loop is fine as long as every thread executes the same number of iterations of that loop. A loop whose trip count depends on threadIdx.x will have some threads exit early and the barrier count will be short by exactly those threads’ arrivals. Loops with data-dependent trip counts must either be made uniform, or the varying part must be predicated out and the loop run to the maximum trip count for all threads.

The barriers that predication cannot replace

The reason predicated fills are used in tiled kernels — described in tail handling and bounds checks — is that the barrier needs every thread in the block to reach it, and an early return removes threads from the block’s accounting permanently. The fill with a neutral value is not a stylistic alternative to a bounds check; it is the only way to keep the barrier well-defined when the tile extends past the end of the buffer.

The same applies to shared-memory allocations. A thread that has returned early has not claimed its share of a static shared array, but static shared memory is allocated per block regardless, so returned threads cost shared memory without using it. This is only a problem for the block’s occupancy, not its correctness.

Warp scope is finer and cheaper

Within a warp, threads can exchange register values directly, without shared memory and without a barrier. The primitive is a shuffle, and its synchronisation requirement is expressed as a member mask: the set of lanes that must participate. __shfl_down_sync(0xffffffff, v, 16) moves lane $k+16$‘s value of $v$ into lane $k$, for every lane in the mask. A reduction across 32 lanes is five of these, halving the offset each time.

The cost profile differs sharply from a block barrier. A shuffle is a register-to-register move with a latency of a few cycles, and it requires no shared memory and no counter. A block barrier requires all warps to arrive, so its cost is that of the slowest warp; if another warp in the block is still executing a long memory operation, every other warp waits at the barrier for the duration of that latency. A tree reduction that uses shared memory plus barriers pays that cost once per level; a warp-level reduction followed by one shared-memory exchange of 32 partial results pays it once in total.

The mask matters. A mask that names lanes which are not executing the shuffle is undefined; a mask that names a subset, used consistently by every participating lane, is a legal way to restrict the operation. Since independent thread scheduling was introduced, lanes of a warp can be at different program counters, so a shuffle no longer implies that the whole warp is converged — the mask is the only statement of which lanes are expected to participate. The classic defect is a loop whose mask is taken from the currently active set: lanes that reach the shuffle at different times compute different masks for the same logical operation, and the exchange silently pairs the wrong lanes.

Named barriers

Hardware provides a small fixed set of barrier resources per block — sixteen numbered barriers in PTX — and the compiler allocates barrier 0 to __syncthreads(). Explicitly naming a barrier allows a block to be split into groups that synchronise among themselves, which turns a block of 512 threads into, for example, two independent halves each running its own pipeline stage. Each named barrier counts arrivals from the threads that execute it, so the group must be defined by a condition that is uniform across the participating threads.

Named barriers are used where a block hosts several independent producer-consumer pairs and the full-block barrier would couple them unnecessarily. The correctness conditions are the same as for the unnamed barrier — every participating thread arrives the same number of times — and the hazard is the same, only harder to see, because the two groups are now different sets of threads and a mismatch between the group’s size and the barrier’s expected count is not detectable by reading the code locally. Named barriers are exposed through inline PTX rather than as a first-class language feature in CUDA C++, which is a fair signal of how rarely they are the right tool.

What a barrier does not do

Three things the barrier does not provide, each of which appears as a bug somewhere.

It does not order accesses across blocks. A thread that writes to global memory before the barrier is not guaranteed to have its write visible to another block after the barrier, because the barrier’s ordering is scoped to the block.

It does not flush or invalidate anything. The caches inside a multiprocessor are coherent with respect to the block’s own stores by construction, so there is nothing to flush; a write to shared memory followed by a barrier followed by a read by another thread in the same block is a simple case of same-location consistency inside one processor.

It does not serve as a scheduling primitive. Placing a barrier to “let the other warps catch up” does not reduce divergence within a warp, does not change the order in which blocks are dispatched, and does not make a slow warp cheaper. It converts a lack of synchronisation into a stall, which is only the right trade when the data actually must be exchanged.

Where the cost lands

The cost of a block barrier is dominated by the number of warps that must arrive and by the variance in their arrival times. A block of 1024 threads is 32 warps; each barrier requires all 32 to reach it, so the barrier effectively serialises the block at the pace of whichever warp is furthest behind. The same computation split across four blocks of 256 threads has four independent barriers, each over 8 warps, and no coupling between them.

That is the design constraint behind small blocks in barrier-heavy kernels. The tiled algorithms that need frequent synchronisation are usually written with several small blocks rather than one large one, even though a large block would hold more shared memory in one place and could compute a larger tile with fewer halo loads. The extra halo traffic is the price paid for a shorter critical path at each barrier, and the exchange rate between them is the central tuning decision in tiled kernels.