System-Scope Atomics and Memory Ordering Across the Bus

Device memory is not coherent with host memory in the way that the caches of a multicore CPU are coherent with each other. On a CPU, a store to a cache line owned by another core triggers a coherence transaction that invalidates or updates the other copy, and the hardware guarantees that all cores eventually agree on the value. Between a GPU and a host CPU connected by PCIe, no such mechanism exists. There is no bus transaction that invalidates a host cache line when a device writes the corresponding device memory, because the device memory is not backing that cache line at all; the two are different physical storage.

That asymmetry is why the memory model has an explicit scope dimension. An atomic operation is atomic with respect to some set of observers, and choosing the scope chooses the set. Getting the scope wrong produces a program that is correct on one platform and silently broken on another, or correct in a test and broken under load, because the failure depends on timing that only manifests when the machine is busy.

What atomicity means without coherence

An atomic read-modify-write on device memory is performed by hardware at the memory controller or in the L2 slice that owns the line. It serializes against other atomics to the same address from other threads on the same device, and it returns the previous value. That much is guaranteed within the device.

The host’s view is different. If the host maps device memory into its address space, which is possible over the bus for some architectures and configurations at substantial latency, a host read of that address goes through the mapping and sees whatever the memory controller holds. There is no cached copy on the host to be stale, because the mapping is uncached. The atomic is visible, but every access costs a bus round trip.

If the buffer instead lives in host memory and the device accesses it through a mapping, the situation inverts: the host’s cached copy can be stale with respect to the device’s writes, and the device’s writes bypass or invalidate host caches only at the granularity the platform implements. A device write to host memory is a DMA write; whether it invalidates the corresponding host cache line depends on whether the platform supports cache-coherent DMA at that granularity. On x86, PCIe writes to host memory are coherent with CPU caches by protocol; on many ARM server configurations they are not unless the interconnect implements coherency for that device, and the driver must manage cache maintenance explicitly.

The practical consequence is that a flag written by a device and polled by the host must be written with an operation whose scope includes the host, and the host must read it with an operation that cannot be hoisted or cached. A plain store followed by a plain load is not a synchronization pattern; it is a data race that happens to work on the platforms where the compiler and the hardware are forgiving.

Scopes and what each one costs

The memory model exposes scopes that correspond to the sets of observers the hardware can actually reach. A block scope covers the threads of one block, which share an L1 cache and a barrier. A device scope covers all threads on one device, which share L2 and the memory controllers. A system scope covers the device, the host, and any peer devices that the platform keeps coherent.

The cost ordering follows the topology. A block-scope atomic is resolved in L1 or in a shared-memory location and costs tens of cycles. A device-scope atomic goes to L2 and costs on the order of a hundred cycles. A system-scope atomic may require the operation to bypass caches entirely, or to be accompanied by a flush, so that every observer sees it; the cost is on the order of hundreds of nanoseconds, dominated by the bus round trip rather than by the arithmetic.

Choosing a wider scope than necessary is therefore expensive, and choosing a narrower one is wrong. The narrow choice fails only when the observer is outside the scope, which is exactly the case a test on a single device will not exercise.

#include <cuda/atomic>

/* Producer on the device publishes a payload, then a flag.
   The release store orders the payload write before the flag write,
   and the system scope makes the flag visible to the host. */
__global__ void produce(int *payload, cuda::atomic<int, cuda::thread_scope_system> *flag)
{
    if (threadIdx.x == 0) {
        payload[0] = 42;
        /* Release: earlier writes become visible before this store does. */
        flag->store(1, cuda::std::memory_order_release);
    }
}

/* Host side. The flag must live in memory the host can reach: managed,
   host-mapped, or a mapped device window. Acquire pairs with the release
   above, so the payload read cannot be reordered before the flag read. */
void consume(int *payload, cuda::atomic<int, cuda::thread_scope_system> *flag)
{
    while (flag->load(cuda::std::memory_order_acquire) == 0) {
        /* spin; on real hardware this should back off or use a
           blocking wait so the core is not starved */
    }
    int v = payload[0];   /* guaranteed to be 42 */
    (void)v;
}

The pair of operations is what matters, not either one alone. A release store without an acquire load on the other side orders nothing: the consumer is free to read the payload before the flag. An acquire load without a release store gives the consumer a guarantee that is not backed by anything the producer promised. The C++ memory model defines these as a synchronizes-with relation between a release operation and an acquire operation that reads the value written by it, and the device implementation is required to honour that relation for the scopes involved.

The trap of a flag in ordinary memory

The most common error is a flag in a buffer that the device writes with an ordinary store and the host polls with an ordinary load. Three things can go wrong, and all three are invisible in a debugger.

The compiler can keep the host’s load in a register across iterations of the polling loop, so the loop never terminates even after the device has written the flag. Declaring the flag volatile prevents that particular transformation but does not provide ordering with respect to the payload. The correct fix is an atomic load with acquire semantics.

The second failure is ordering. Without a release on the producer side, the payload store and the flag store may become visible in the opposite order. A consumer that sees the flag set may read the payload before the producer’s write has landed. On a device this can happen because the two stores go to different memory partitions and complete independently; the flag is not a barrier.

The third is granularity. If the flag shares a cache line with the payload, a consumer polling the flag causes the line to be fetched repeatedly, and on a platform where the device’s writes to that line require invalidation, the polling traffic can interfere with the producer’s progress. Flags should sit on their own cache line, padded to 64 bytes or 128 bytes depending on the architecture’s line size, so polling does not contend with payload traffic.

Where the model interacts with the stream abstraction

Stream ordering and the memory model answer different questions, and confusing them produces code that is correct in one dimension and broken in the other. A stream guarantees that a kernel submitted later begins after an earlier one completes, and that a synchronization of the stream makes prior writes visible to the host. It says nothing about two threads communicating through a flag while both are running, and nothing about a host thread reading a buffer while a kernel is still executing.

A host thread polling a flag written by a running kernel is therefore outside the stream model entirely. The stream will not be synchronized until the kernel completes, and the kernel will not complete until it finishes, so the only way to observe intermediate progress is through the memory model, with system-scope operations and a mapping the host can read. The stream’s synchronization call remains necessary afterwards, to establish that the kernel has finished and to flush anything the memory model left pending.

The reverse direction has the same shape. A kernel polling a flag that the host sets requires the host’s store to be visible to the device. On a platform where host writes to device memory are not coherent, the host must use a store with release semantics and a scope that includes the device, and the mapping must be one the device reads without caching. Where device memory is mapped into the host’s address space, the access is uncached and the cost is a bus round trip per poll, which makes tight polling loops on the host expensive.

Practical rules that hold across platforms

Use the narrowest scope that includes every observer. Block scope for intra-block cooperation, device scope for cross-block coordination on one device, system scope only when the host or a peer device participates.

Pair every release with an acquire. A release store with no matching acquire load provides no guarantee to the reader; an acquire load with no matching release provides none to the writer.

Pad synchronization variables to a cache line. A flag sharing a line with frequently written data turns every poll into traffic that competes with the writes it is waiting for.

Prefer the stream and event machinery for coarse coordination. Atomics are for fine-grained communication where a kernel must observe another’s progress without ending, and for lock-free data structures. Using them to sequence whole phases of a pipeline replaces a hardware queue with a software spin loop and usually loses.

Finally, treat any pattern that relies on the absence of caching as platform-specific. The correctness argument has to name the scope, the ordering, and the mapping attribute that make it valid. Code that works because the test machine happens to be coherent is not a memory model; it is a coincidence. The mapping attributes that determine whether an access is cached are described in host-mapped memory over PCIe, and the queue-level ordering that atomics are usually a poor substitute for is covered in stream ordering and submission semantics. On systems where the kernel maintains page tables for both processors, the coherency question is partly answered by the mirroring described in page table mirroring between CPU and device, though mirroring establishes translation agreement, not cache coherence.