100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryPerformance
CC 7.5

What is atomic contention in CUDA?

Many threads hitting the same address with atomics, which serializes them no matter how many threads you launched.

What you pay for is the address, not the instruction. An atomicAdd reads a word, adds to it and writes it back so the three steps "appear to execute in a single step", and that promise is per address: two atomics on two words go through together, two on one word take turns. A million threads adding into one float is a queue a million deep with a GPU attached to it, and every thread you add makes it longer. The same four-line kernel is fast over a million counters and slow over one.

Two things an atomic is not, and the failures look unrelated. It is not a barrier: "legacy atomic functions only ensure atomicity and do not introduce synchronization points (fences)", so a thread that must see what another thread wrote still needs __syncthreads() from day 14. And it is not an order. These carry "a memory ordering of cuda::std::memory_order_relaxed", so the hardware performs your n additions in whatever sequence suits it and a different one next run. For integers that costs nothing. For floats it costs the last bits, because addition is not associative: on values that are not exactly representable, day 26's global-atomic sum returned 20 distinct totals across 20 runs where the tree returned 3. That is floating point determinism.

The folklore is that 32 lanes queue up, so an atomic costs 32 plain adds. Treat that as an upper bound, because the compiler gets there first: nvcc "performs warp aggregation for atomics automatically in many cases", electing one thread to add an increment the warp already totalled among itself. Read the case NVIDIA documents before you claim it: a counter, every lane adding the same 1 to one address, which a population count folds into a single add. Summing an array gives every lane a different value, so folding needs a real warp shuffle reduction of the kind day 23 writes by hand. Maxwell also replaced Kepler's "lock/update/unlock pattern that could be expensive in the case of high contention" with native shared-memory atomics, for 32-bit integers, not a float add. Measure your own card.

The fix is to move the contention down a level, and there are two levels to move it to. Privatization gives each block its own accumulator in shared memory and sends one atomic to the global one, which changes which memory serializes rather than whether it does. A tree inside the block removes the serialization and still sends one global atomic per block, so the only difference between those two is how a block folds its 256 values into one. The ranking that falls out does not hold across sizes.

Measured

Day 26 sums one buffer three ways at five sizes, 10 timed runs after 3 warm-ups, each launch getting a fresh accumulator. All three read every element once with consecutive threads on consecutive addresses, so the coalesced read is identical and only the accumulation moves. It is a float add throughout, on a compute capability 7.5 card. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75.

n global atomic shared privatised tree fastest
4,096 0.0175 ms 0.0300 ms 0.0047 ms tree
65,536 0.2333 ms 0.0774 ms 0.0097 ms tree
262,144 0.9229 ms 0.2370 ms 0.0259 ms tree
1,048,576 3.6827 ms 0.8522 ms 0.0918 ms tree
4,194,304 13.1734 ms 1.5241 ms 0.1666 ms tree

At the largest size that is 1.3, 11.0 and 100.7 GB/s counting the four bytes per element each kernel reads, and the tree finishes 79 times faster than the naive version. The gap widens with n, which is the signature of contention: the tree's work per element is fixed while the queue at the global accumulator grows with the launch. Every row returned the same answer bit for bit against a Kahan double reference, so the table prices contention alone.

Read the first row before you generalise. At 4,096 elements the privatised version is the slowest of the three, 0.0300 ms against 0.0175 ms: sixteen blocks on one address is barely a queue, and privatising spends two barriers to remove a cost that was not there.

The data decides how contended a kernel is, not only the code. Day 29 puts 67,109,475 bytes into 256 bins: uniform bytes give the busiest bin 0.39 percent of the updates and the global-atomic kernel takes 21.703 ms, one repeated byte gives it 100.00 percent and 55.400 ms. Privatising costs 0.759 ms and 2.386 ms on those two inputs.

Diagram

An original SVG, atomic-contention-three-levels. Three bands, each with 256 thread boxes on the left, a block's shared memory in the middle and one global accumulator on the right. Band one runs 256 arrows straight past shared memory into the global box. Band two converges 256 arrows on one shared box with a single arrow leaving it. Band three folds the 256 through eight halving rows inside shared memory, with one arrow out.

Alt text: "Three routes from one block of 256 threads to a total. Two hundred and fifty-six adds serialized on one global address, then the same count moved to a shared address and one global add, then no serialized adds at all and eight barrier steps."

Code

From code/day26-atomics/atomics.cu. The naive kernel's whole body is atomicAdd(total, in[i]); under a bounds check. The tree replaces it with this, and the point is the last two lines: the number of global atomics falls from one per element to one per block.

    for (unsigned int half = kThreadsPerBlock / 2; half > 0; half /= 2) {
        if (tid < half) {
            tile[tid] += tile[tid + half];
        }
        __syncthreads();
    }

    if (tid == 0) {
        atomicAdd(total, tile[0]);
    }

The privatised version sends the same one global atomic per block. It differs only in folding those 256 values with a shared atomicAdd instead of the loop above, which is why the two columns in the table diverge at all.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. The numbers came off the verification node on 2026-08-30, and this entry stays a draft until a named author and a different named reviewer sign it.