Atomics and Contention in Counters and Aggregates

A global atomic increment is the simplest cross-block communication primitive and the one most often used where it should not be. Its semantics are strong and easy to reason about: the read-modify-write completes as a single indivisible step, so no increment is lost and the counter’s final value is the number of increments. Its performance profile is the opposite of its simplicity. All the increments in a grid target one location, that location lives in one cache line, and the memory system must serialise the updates to that line.

What the hardware does with one contended address

An atomic to global memory for a block that is not resident in the executing multiprocessor’s L1 is performed at the L2 slice that owns the line. The modern path does this without a read-modify-write round trip to the requester: the arithmetic is performed inside the L2, so the line is not handed back and forth, only the operation. This makes a single atomic far cheaper than a load-bracketed-store sequence, but it does not make it free, and it does not make it parallel.

The throughput limit is per line. Every update to the same address must be ordered against every other, and the L2 processes them one at a time. Thirty-two lanes of a warp issuing an atomic to the same address are therefore served by a serialised sequence of updates, not by one, and the instruction occupies the warp for the duration of the whole sequence. Implementations do detect the special case where every active lane targets the same address and can collapse it into one operation, but that reduction applies to the fully populated, fully uniform warp and to nothing else: a warp with 30 active lanes, or with lanes spread over three addresses in one line, pays the per-operation cost for all of them.

Two consequences follow. First, an atomic that aggregates a value across the whole grid is a serialisation point whose cost scales with the number of warps, not a per-thread operation with constant cost. Second, the same atomic issued by different blocks contends at the same L2 slice, so the cost does not distribute itself across the memory system the way ordinary loads and stores do.

Aggregating before the atomic

The remedy is to reduce contention by reducing the number of atomic operations, not by making each cheaper. The standard ladder has three rungs.

Within a warp, a shuffle-based tree computes the warp’s contribution and a single lane issues the atomic. A warp of 32 threads doing this turns 32 atomics into one — a 32-fold reduction in traffic with five shuffle-and-add steps. The reduction is over the count of atomic operations, which is the quantity that matters, and the shuffle cost is the same as in any tree reduction.

Within a block, the per-warp results go into a small shared array, a barrier, and one warp reduces them. The block then performs one atomic. For a block of 1024 threads that is 1 atomic per 32 warps, down from 1024.

Across the grid, the remaining $B$ atomics — one per block — can still contend. The final rung is to replace the single counter with an array of counters, one per slice, and sum them afterwards. A thread hashes its block index to a slice, and each slice is contended by only $B / S$ blocks. The slices must be far enough apart to land on different L2 slices and, more importantly, on different cache lines: two counters in the same line contend for the same line even though they are different addresses. Padding the array so that each counter occupies a full line, 128 bytes on the common architectures for global accesses, is necessary, not optional.

// One padded counter per slice; each occupies its own 128-byte line.
struct alignas(128) PaddedCount { unsigned long long v; };

__global__ void count_matches(const int* data, size_t n, int key,
                              PaddedCount* slices, int nslices) {
    unsigned long long local = 0;
    for (size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += (size_t)blockDim.x * gridDim.x) {
        local += (data[i] == key);
    }
    for (int off = 16; off > 0; off >>= 1)
        local += __shfl_down_sync(0xffffffffu, local, off);
    if ((threadIdx.x & 31) == 0) {
        // Blocks are spread across slices; one atomic per warp per slice.
        unsigned long long* p = &slices[blockIdx.x % nslices].v;
        atomicAdd(p, local);
    }
}

The sum over slices is done on the host or in a small follow-up kernel. The count is exact — integer addition is associative — so the aggregation changes nothing about the result, only the cost.

Ordering, and what the atomic does and does not guarantee

An atomic operation is indivisible and linearly ordered with other atomics to the same location. That is all. It does not order the thread’s other memory operations against it, and it does not make a write to a different address visible to another block.

The common pattern that is broken by this: a block writes a result to a buffer, then increments a counter to announce that the result is ready. If the write is an ordinary store and the increment is a relaxed atomic, another block that reads the counter and then reads the buffer may see the counter’s new value before the buffer’s contents have become visible. The buffer read then returns stale data, and the bug appears as an intermittently wrong value under load.

The fix is a release on the publishing side and an acquire on the observing side. The thread that writes the buffer must perform its announcement with release semantics, so that the buffer writes cannot become visible after the announcement; the thread that observes the announcement must perform that read with acquire semantics, so that its subsequent read of the buffer cannot be reordered before it. In CUDA these are expressed through the order and scope of the atomic operation, or through an explicit fence between the write and the announcement. The same hazard appears in the publication of scan descriptors discussed in decoupled look-back, where the remedy used is to make the state and the payload a single atomic word so that no ordering is needed at all.

Scope is a separate axis from order. A device-scope operation is ordered with respect to other accesses from the same device; a system-scope operation is ordered with respect to host memory and peer devices. The legacy atomic functions on device memory provide no ordering by themselves, which is why an explicit fence or an atomic with release and acquire semantics is required wherever an atomic is used to publish data. Choosing the narrowest scope that covers the participants is legitimate, but it must be a deliberate decision: the failure mode when it is wrong is data visible later than expected, which does not reproduce reliably.

Floating-point atomics and non-determinism

atomicAdd on floats exists and is atomic in the sense that it performs an indivisible read-modify-write. The addition it performs is not associative, so the final sum depends on the order in which the increments arrive. The order depends on scheduling, so two runs over the same input can differ in the last bits. This is a property of the algorithm, not a defect: any parallel sum of floats has an order, and the atomic simply makes the order unpredictable.

Where that matters, the options are to aggregate deterministically before the atomic, so that the number of atomic operations and their order are fixed, or to accumulate in a format that is associative — integer fixed-point with a known scale, or a compensated scheme that carries an error term. The first is usually enough: a block-level reduction followed by one atomic per block has at most $B$ atomic arrivals, and if $B$ is small the non-determinism is confined to a few orderings that can be enumerated.

Atomic contention as a symptom

A kernel whose performance collapses when a counter is added has found a structural problem, not a micro-optimisation target. The usual diagnoses are:

The counter is on the critical path of every thread, so its rate is bounded by the atomic throughput rather than by the work. This is the case for a global work queue where each thread pops one item: the queue head is a single contended location, and the pop rate is the atomic rate. The fix is to hand out work in batches — each thread pops $k$ items — so the atomic rate falls by $k$.

The counter’s line is being fought over by blocks that also need the data in that line. False sharing between a counter and a result buffer is easy to create and hard to see; placing counters in a separately allocated, line-aligned region removes it.

The atomic is being used where an ordinary store would do. A counter written by exactly one block needs no atomic at all. A flag set once needs a single store with release semantics, not a read-modify-write. Every read-modify-write on a location that has one writer is wasted serialisation.

Where atomics are still the right answer

Contention is only a problem when the rate is high. An atomic used once per block, or once per grid, is a few dozen operations and costs nothing measurable; the code is simpler than a descriptor protocol and the ordering is easier to state. The decision boundary is the number of operations per unit of work, not the presence of the atomic itself.

The patterns that genuinely need atomics are the ones where the final order is inherently dynamic: a compaction that appends results to a buffer in whatever order the threads finish, a histogram whose bins are spread by hash, a work counter that hands out ranges. All of these can be improved by batching and by spreading across slices, and all of them remain correct only if the aggregation preserves the semantics the algorithm needs. In the compaction case that means the result order is not stable across runs, which is the same non-determinism as the atomic float sum and has the same remedies.