Atomic Contention and the L2 as a Serialization Point

An atomic operation is the memory access whose cost depends most directly on how many other threads are doing it at the same time. Everything else in the hierarchy is per-access: a load costs sectors, a scratchpad access costs passes, a DRAM access costs a burst and sometimes a row sequence. An atomic on a contended location costs a queue, and the queue is at a fixed physical place in the machine. That place is the slice of the last-level cache that owns the address, which is also the point at which the device’s writes become visible. Understanding contention is therefore mostly a matter of understanding where one address lives and how much work can arrive there at once.

The operation executes where the address lives

A global atomic is not performed by the SM that issued it. The request travels to the cache slice responsible for the address, and the read-modify-write happens against the bank in that slice, with the slice’s own arbitration deciding the order in which the operations from different SMs are applied. This is the same placement rule that applies to ordinary loads and stores, and the reason it must be so is coherence: the modification has to happen at the point every other thread’s read is checked against, otherwise two SMs could each hold a copy of the old value.

The consequence is a locality property that holds for atomics and not for ordinary reads. Two threads incrementing the same counter in two different blocks contend no matter how the blocks are scheduled, because the operation is serialized at one address. Two threads incrementing two counters two sectors apart may be served by the same slice, and one slice can only apply operations at its bank’s rate; the operations are not to the same location but the resources overlap. Spreading atomic targets across the address space spreads them across slices, which is why a histogram with a bucket array is fundamentally more parallel than a sum into one accumulator.

Atomics in shared memory behave differently, and this is worth separating because the difference is large. Such an operation executes on the SM, in the same array and under the same bank rules as any other scratchpad access. It is serialized with other atomics to the same word, but the round trip is on-chip rather than across the device interconnect, so the queue drains much faster. Contention remains the constraint; its price changes by an order of magnitude.

Why a hot counter stops scaling

The bank in a cache slice applies read-modify-writes at some rate, and the round trip between an SM and the slice has a latency of a few hundred cycles. A single address being incremented by every thread of a large grid therefore becomes a pipeline in which the slice’s rate, not the number of SMs issuing, sets the throughput. Adding more threads makes the queue longer rather than the work faster, and the throughput available per thread falls in proportion to the number of contenders.

The scaling limit is per address, not per atomic instruction. A reduction loop in which each atomic targets a different element of an array scales with the array length until the addresses collide in the hash, which is why a bucket array in a histogram wants a length whose buckets do not resonate with the slice mapping, the same consideration that governs the slice mapping. A counter incremented by every thread has no such escape.

There is a second-order cost in read-modify-write loops. When the update is expressed as a compare-and-swap loop, a failed comparison means the loop repeats, and contention increases the failure rate. Each retry is another full round trip, so the effective rate per successful update degrades faster than linearly with the number of contenders. The failure is not livelock — a lock-free loop guarantees that some thread makes progress — but the wasted bandwidth is real.

Removing contention before optimizing it

The standard reduction rung is to move the accumulation into the block before it touches global memory. Each block reduces its own partial sums in shared memory, using the scratchpad hierarchy described in the tiling discussion, and then one thread per block performs a single global atomic. The number of global atomic operations falls from one per thread to one per block, which for a block of a few hundred threads is two orders of magnitude, and the remaining operations still contend but far less often.

The scheme generalizes. Any associative combination can be applied hierarchically, with each level reducing the number of contending operations by the width of the level. A per-thread partial, a tree reduction across the lanes of a warp, a warp-level result accumulated into a shared array, and one final global atomic per block give four levels, and the shared-memory levels are cheap because their atomics are local.

Reductions that need not be atomic at all are the next step. If the shape of the data permits a deterministic tree — a fixed sequence of pairwise combinations — then no atomic instruction is required, and the result is reproducible. The cost is that the tree has a fixed shape, so the parallelism of the reduction is bounded by the number of nodes at its widest level, and the implementation must handle a size that is not a power of two.

// Block-level reduction: shared memory for the tree, one global atomic at the end.
// Launched with 256 threads per block; a power-of-two block size keeps the
// halving loop simple.
__global__ void sum(const float *in, float *out, int n) {
    __shared__ float partial[256];
    int t = threadIdx.x;
    float acc = 0.0f;
    for (int i = blockIdx.x * blockDim.x + t; i < n; i += gridDim.x * blockDim.x)
        acc += in[i];                    // coalesced: lanes differ by 1
    partial[t] = acc;
    __syncthreads();
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (t < s) partial[t] += partial[t + s];
        __syncthreads();                 // the tree needs a barrier per level
    }
    if (t == 0) atomicAdd(out, partial[0]);   // one atomic per block, not per thread
}

The tree uses a barrier per level, which costs synchronization but keeps the accumulation in registers and shared memory. Whether the last step should be an atomic or a second kernel that reduces one partial per block is a separate decision: an atomic avoids a launch but makes the result order-dependent, while a second kernel gives a deterministic ordering at the cost of another pass over a very small array.

The form without a return value

An atomic operation that returns the previous value must round-trip that value back to the issuing thread, which occupies interconnect bandwidth and forces the thread to wait for it. Many uses do not need it. When a result is only accumulated and never inspected by the issuing thread, the operation can be issued in a form that discards the old value, and the wording of the instruction set distinguishes the two: a reduction operation with no return path.

The distinction matters for throughput in two ways. It removes the response traffic, which on a contended address is a significant fraction of the cost, and it allows the hardware to combine or reorder operations to the same location more freely, since there is no ordering relationship implied by returning a value to a specific thread. The compiler selects this form automatically when the return value is unused, so the gain is available without naming it, but it is lost if a kernel stores the returned old value into a variable that is never read again in a way the compiler fails to see.

Relaxed by default, ordered on request

Atomic operations in the programming model are relaxed by default: they guarantee that the modification itself is indivisible, not that operations before or after it are visible in a particular order. A kernel that uses an atomic as a synchronization point and then reads a data structure written by another thread needs an explicit ordering mechanism, either a fence or an acquire-release operation with a declared scope.

The scope is the other half of the question. An operation ordered at block scope does not order anything for another block, and an operation ordered at device scope does not constrain the host. A kernel that uses block-scoped ordering to publish data to a peer block is relying on a guarantee the hardware does not make, and the failure appears as a read of stale data under a workload the developer cannot control.

The practical shape of all of this is a hierarchy of decreasing contention: per-lane values, a shared-memory tree, one atomic per block, and only when the target is genuinely a single global location does the contention return. A design that treats atomics as ordinary memory operations and then finds the counter is the bottleneck is usually fixed by the multi-level reduction, not by a faster atomic instruction, because the instruction was never the constraint. The address was.