Atomics, Fences, and Memory Scopes in Device Code

An atomic operation on a GPU is not a locking operation. It is a read-modify-write on one address, executed by the memory subsystem, with an ordering parameter that says how much of the surrounding memory traffic must be ordered around it. The distinction is the source of most incorrect lock-free device code: atomicAdd guarantees that no update is lost, and says nothing about whether the data the updated value refers to is visible.

Three properties that vary independently

An atomic instruction has three properties that programmers routinely conflate.

Atomicity is per operation: the read and write of one address happen as one indivisible step. Two blocks incrementing the same counter produce the correct count. This holds even for the relaxed variants.

Ordering is per program: whether the atomic acts as a barrier with respect to the surrounding loads and stores in the same thread. A relaxed atomic may be reordered with any other memory operation; a release atomic prevents preceding operations from moving after it; an acquire atomic prevents succeeding operations from moving before it.

Scope is per population: which other threads must observe the effects. Device scope covers all threads on the GPU; block scope covers the block; system scope extends to the host and to peer devices. A block-scope atomic gives no guarantee at all about threads in another block, even if the operation was atomic.

A fourth property hides in the API’s naming rather than its semantics: visibility point. A value written to global memory by one thread becomes readable by another only after the write has reached a level of the memory hierarchy both can access. A fence is the instruction that forces this; an atomic with release ordering implies one.

Why relaxed atomics do not order anything else

The canonical pattern is a flag carried by an atomic and payload carried by ordinary stores. The producer needs a release-like ordering and the consumer needs an acquire-like one, and the second is the one usually left out:

__device__ int flag = 0;
__device__ float payload;

__global__ void produce(float v) {
    payload = v;
    __threadfence();              // payload store visible before the flag store
    atomicExch(&flag, 1);
}

__global__ void consume(float *out) {
    while (atomicAdd(&flag, 0) == 0) { }
    __threadfence();              // without this, the load may hoist above the loop
    *out = payload;
}

The producer’s __threadfence() orders the payload store before the flag store, so no thread that observes the flag can observe a stale payload. Delete the fence in consume and the guarantee disappears: a relaxed atomic read carries no ordering, so the compiler or the hardware may hoist the load of payload above the loop exit. An atomic with acquire semantics achieves the same thing at lower cost where the architecture supports it, because the ordering can be fused into the load rather than issued as a separate barrier.

The reason such a program often appears to work is that the flag load sits in the loop condition, where it has to stay, so the defect appears only when the compiler moves the load of payload above the loop, which relaxed semantics permit, or when the hardware reorders the two. Both require an optimiser or a memory system willing to use the freedom the model grants, so a short test run rarely shows the bug.

Fences, and what they actually flush

__threadfence() orders the calling thread’s memory operations as observed by the whole device: writes before the fence become visible to every thread on the device before any write after it. __threadfence_block() is the weaker version, ordering only within the block, and it does not guarantee that a write has left the multiprocessor’s L1.

That distinction is not academic. A write that hits a write-back L1 cache is visible to the writing thread’s block and to nothing else until it is evicted or flushed. A block-scope fence does not force the flush; a device-scope fence must. Where the L1 is write-through or not coherent across multiprocessors, a device fence is the transition that makes cross-block communication work at all.

The cost is the point. A device-scope fence waits for outstanding writes to reach a coherence point accessible to every multiprocessor, which is potentially the L2. A block-scope fence is cheaper, and an acquire or release attached to an atomic may be cheaper still, because the memory subsystem can fuse the ordering into the atomic’s coherence transaction rather than treating it as a separate barrier.

The general rule: use the narrowest scope that covers the threads that must observe the effect. Two blocks coordinating a partial sum need device scope because they cannot see each other’s caches otherwise. Two warps inside one block coordinating through shared memory need block scope only.

Scopes and the difference between shared and global

On-chip scratch memory is not coherent across blocks at all. It is private to one multiprocessor by construction. An atomic on a shared-memory address is therefore only meaningful within the block, and the implementation is free to reject a device-scope qualifier on it or to treat it as block scope.

Global memory is the only medium across blocks, and its coherence point is the L2, which is shared by all multiprocessors on the chip but not by another chip in a multi-GPU system. System scope names the wider population, which is what a thread doing cross-device communication or communicating with the host through mapped memory must use. The cost of system scope is the flush of L2 lines that are visible to the peer or host, and it is higher than the device-scope equivalent. Which scope a device actually offers is a property the API reports, so it belongs with the other queries in Capability Negotiation Across Architectures.

CUDA spells it with _block and _system suffixes on the atomic functions, the unsuffixed name meaning device scope, and with a memory-order argument on the cuda::atomic types. OpenCL 2.0 spells it as an explicit memory_order argument together with a memory_scope argument on atomic_* functions, which is the clearer formulation because the two dimensions are separate parameters rather than an enumerated product.

Practical classes of lock-free kernel

Three patterns recur, each with a scope requirement that follows from its structure.

A device-wide counter incremented by every work-item needs device scope and relaxed ordering for the increment itself, since no other memory is being ordered, plus a release/acquire pair only at the point where the counter’s value is used to index into a result array. Wrapping the increment in a stronger ordering than necessary is the common over-specification.

A block-level reduction through shared memory needs block-scope atomics or a block barrier, never a device fence. The barrier is also the cheaper primitive where a full exchange is not required, and its cost is the wait for every thread in the group to arrive.

A queue built from a single head index published to other blocks needs a release store on the head and acquire loads on the consumers, both at device scope. A single-producer ring with a monotonically increasing index is the version that avoids a compare-exchange loop entirely, because the producer is the only writer and atomicAdd returns a unique slot.

Where the model stops

The ordering guarantees are only as strong as the implementation’s, and the specification is deliberately permissive about scheduling. Two work-items in the same group are not guaranteed to make simultaneous progress unless they reach a barrier, which means a spin-wait between them can deadlock if the hardware schedules groups in a way that never runs the producer. Forward-progress guarantees for spinning code exist at block granularity, since a resident block will eventually run, but not within a block unless the code meets the conditions for independent thread scheduling.

This is the deepest asymmetry between device code and host code, and it does not surface in the atomic primitives themselves. It surfaces in algorithms that assume a spinning consumer will eventually be scheduled, which is routine on a CPU and unsafe on a machine whose warps are interleaved by a hardware scheduling policy. The safe constructions are the ones where every participant reaches a barrier, or where the spin is bounded and failure is handled. The same asymmetry decides which ordering guarantees a program may rely on at the launch boundary, where the queue semantics described in Command Queues, Events, and Cross-Stream Ordering take over.