Barriers, atomics, and memory order
Choose a barrier, atomic operation, fence, or separate kernel from the participant set, shared data, and required visibility order
Before you begin
Follow a linked prerequisite when you cannot explain it from memory. Background skills without links are stated as abilities rather than hidden terminology.
- Ability to identify two concurrent accesses to one memory location
- Chapter 4 definitions of warp, block, shared memory, and device memory
Suppose 256 threads load one tile into shared memory and then compute with values loaded by other threads. A thread that finishes its own load cannot immediately begin the computation. Another thread may not have written the required element yet. The program needs a phase boundary after the cooperative load and before any cross-thread read.
Now consider a histogram. Several threads may increment the same bin. A barrier does not make those increments safe because the read-modify-write operations can still overlap. The program needs an atomic operation, or it must first assign updates to separate locations and combine them later. The tile and histogram examples require different guarantees.
Synchronization has at least three dimensions: participation, atomicity, and memory ordering. Participation identifies which threads must wait or communicate. Atomicity prevents an operation on one location from being torn apart by a competing update. Memory ordering defines when writes to surrounding data become visible. CUDA provides different primitives because no single operation supplies every guarantee at every scope.
Barriers coordinate a known participant group
A block barrier such as __syncthreads() waits until the required threads in the block reach the barrier. The barrier also provides the documented ordering needed for the common pattern in which threads write shared memory before the barrier and read one another's data after the barrier. The resulting phase ordering makes the cooperative tile correct.
The barrier must be reachable by its participant set. If some threads take a branch containing __syncthreads() while other threads in the same block skip it, the waiting threads may never observe the required arrivals. Structure the bounds checks so threads that have no valid input still participate in the barrier when later phases require the complete block.
A warp-scoped operation can be sufficient when only lanes in one warp communicate. Warp scope avoids waiting for unrelated warps, but warp scope is correct only if no dependency crosses the warp. A device-wide phase is normally expressed as a separate kernel launch unless the program deliberately uses cooperative launch features with their additional constraints.
Atomics protect one shared update
An atomic read-modify-write ensures that competing operations on the target location are applied without lost updates according to the specified operation and memory semantics. One atomic operation does not make an entire data structure atomic. Two atomics on two fields can still be observed between updates unless a larger protocol prevents the intermediate observation.
Contention determines cost. If every thread performs atomicAdd on one global counter, the operations must be coordinated at that location and may queue. Hierarchical aggregation reduces the number of global operations: combine values within each warp, combine warp results within a block, and have one thread update the global counter. The mathematical equivalence must account for floating-point order, overflow, and any ordering requirements.
A barrier and an atomic can appear in the same algorithm. In a block-local histogram, threads may atomically update shared bins, wait at a barrier so all updates complete, and then cooperatively write the bins to global memory. Each primitive addresses a distinct part of the dependency.
Publishing data requires ordering as well as a signal
A common protocol writes a result record and then increments a completion counter. The counter update may be atomic, but that fact alone does not define when a consumer may observe the ordinary stores that wrote the record. The required invariant is stronger: once the completion signal is visible, every field of the corresponding record must also be visible.
Express this invariant with matching release and acquire behavior at the required scope, or with the CUDA primitives documented for the participating agents. A device consumer and a host consumer may require different scope. A fence orders memory effects but does not by itself make threads wait. A barrier makes participants wait and supplies specific ordering guarantees, but it does not combine conflicting writes.
Concurrency failures are schedule-dependent. A test can pass repeatedly because one observed schedule happens to make stores visible in the expected order. Stress tests, Compute Sanitizer race and synchronization checks, and a written invariant are stronger evidence. Performance tuning must begin only after the synchronization contract is correct.
Trace one block-scoped release-acquire protocol
The `producer_consumer_demo` CUDA kernel in the Reference summary launches one block of 64 threads. Thread 0 is the named producer, thread 32 is the named consumer, and both access one record in that block's shared memory. A `cuda::atomic_ref<int, cuda::thread_scope_block>` wraps the shared ready flag. Block scope is sufficient because both participating threads belong to the same block; the scope would not include a consumer in another block.
The producer writes `slot.id` and `slot.score` with ordinary shared-memory stores, then performs `ready.store(1, cuda::memory_order_release)`. The consumer repeats `ready.load(cuda::memory_order_acquire)` until the load reads 1, and only then reads the record. When the acquire load reads the value written by the release store, the release-acquire relationship makes the producer's earlier payload writes visible to the consumer.
Replacing both flag operations with relaxed ordering preserves atomic access to the flag but removes the ordering that publishes the separate payload locations. The consumer can then observe a ready value without a guarantee that it observes the matching record. Moving the consumer to another block while keeping block scope is a second failure: the atomic scopes do not include both participants, so the accesses do not form the required synchronization relationship. Use device scope for a valid cross-block memory-order relation, while also designing a launch and progress scheme that cannot strand a producer behind spinning consumers.
Worked example: Publishing one shared-memory record inside a CUDA block
The `producer_consumer_demo<<<1, 64>>>` kernel assigns thread 0 as producer and thread 32 as consumer. The producer fills one shared-memory record, and the consumer copies the record to global output only after observing a shared ready flag.
- Initialize the ordinary shared integer `ready_storage` to zero in thread 0, then execute `__syncthreads()` so both warps begin only after the flag storage is initialized.
- Construct `cuda::atomic_ref<int, cuda::thread_scope_block>` over that shared integer in the producer and consumer. The block scope names the exact participant set: CUDA threads in this one block.
- Have `produce_record` write both payload fields before it performs the release store of 1 to the flag. The release operation publishes the writes sequenced before it.
- Have `consume_record` wait with acquire loads. After one acquire load reads 1 from the producer's release store, read the two payload fields and compare them with the expected values.
- Run the deliberately incorrect relaxed-order variant under stress. A passing run does not repair the missing ordering guarantee; a stale payload is permitted because atomicity of the ready flag does not order ordinary accesses to the record.
- Keep the consumer in the same block for this example. A consumer in another block requires a scope that includes both blocks, such as device scope, plus an execution design that guarantees the producer can run while consumers wait.
Conclusion: The block-scoped release store and matching acquire load connect the ready signal to the ordinary payload stores. Relaxed flag access or a scope that excludes one participant breaks that connection even though the flag access itself remains atomic.
Common errors
- Using volatile as a general synchronization mechanism. Compiler access rules do not replace atomicity and memory ordering.
- Placing a block barrier inside non-uniform control flow. The barrier's participant set must be valid on every path.
- Choosing global atomics before considering local aggregation. Correctness may be simple, but contention can dominate the kernel.
Reference summary
A barrier coordinates a participant group and provides documented memory ordering for the applicable scope. An atomic protects one read-modify-write operation. A fence orders memory visibility without making participants wait. Select the primitive from the dependency invariant rather than from a general preference for fewer synchronization operations.
A publication protocol often needs two actions. The producer writes the payload, then publishes a signal with suitable release behavior. The consumer observes the signal with matching acquire behavior before reading the payload. An atomic counter update alone does not define visibility for unrelated ordinary stores.
#include <cuda/atomic>
struct Record { int id; float score; };
__device__ void produce_record(Record* slot, int* ready_storage) {
cuda::atomic_ref<int, cuda::thread_scope_block> ready(*ready_storage);
slot->id = 42;
slot->score = 3.5f;
ready.store(1, cuda::memory_order_release);
}
__device__ void consume_record(const Record* slot, int* ready_storage,
Record* observed, int* error) {
cuda::atomic_ref<int, cuda::thread_scope_block> ready(*ready_storage);
while (ready.load(cuda::memory_order_acquire) == 0) { }
Record value = *slot;
observed[0] = value;
error[0] = value.id != 42 || value.score != 3.5f;
}
__global__ void producer_consumer_demo(Record* observed, int* error) {
__shared__ Record slot;
__shared__ int ready_storage;
if (threadIdx.x == 0) ready_storage = 0;
__syncthreads();
if (threadIdx.x == 0) produce_record(&slot, &ready_storage);
if (threadIdx.x == 32) consume_record(&slot, &ready_storage,
observed, error);
}
// Launch exactly one block with at least 33 threads:
// producer_consumer_demo<<<1, 64>>>(observed, error);- Use the narrowest correct scope: lane exchange, warp, block, device, or system.
- Every thread expected at a block barrier must reach it; divergent barriers can deadlock or become undefined.
- Aggregate updates locally before global atomics when contention is high.
- Separate algorithmic phases into kernels when a device-wide barrier is required and cooperative launch is not appropriate.
- Run Compute Sanitizer race and synchronization checks before performance tuning.
Knowledge check
Why is an atomic increment insufficient if a consumer must also observe a producer’s non-atomic data writes?
Show the answer guide
- An atomic increment makes the counter update indivisible, but ordinary payload stores occupy different memory locations and are not automatically ordered by a relaxed atomic update.
- The producer must publish the flag with release semantics after writing the payload, and the consumer must observe that value with acquire semantics before reading the payload.
- The atomic scope must include both participants: block scope works for the named same-block example, while an excluded consumer creates no valid synchronization relationship.
Run the release-acquire publication example
Implement and run `producer_consumer_demo<<<1,64>>>`, then create one relaxed-order variant and document why a passing relaxed run is not a correctness proof. Keep producer and consumer in separate warps of the same block.
Evidence to keep: Source for both variants, launch configuration, correctness output over a stress loop, Compute Sanitizer results where available, and a written scope-and-order invariant.
Synchronization within a kernel is only part of the dependency graph. Streams and events express ordering across kernels, copies, and host orchestration.
Sources and further reading
The lesson explanations synthesize primary papers and official references. Follow the originals for proofs, exact APIs, architecture details, and version-sensitive behavior.