CUDA atomicAdd and what contention costs
Someone on r/CUDA posts a privatized histogram kernel that is already the right shape, and asks why it needs a barrier in the middle:
"I dont know why I need to use syncthreads between these atomic function, In my knowledge that atomic is also wait all the threads execute one by one so why I need to use syncthreads, but if I don't use it, it'll got data race"
https://www.reddit.com/r/CUDA/comments/15226r0/atomic_function_and_syncthreads/ (posted 2023-07-17, checked 2026-08-30)
The question treats an atomic as both an order and a barrier. An atomic does not make threads take turns or wait for each other. It serializes accesses to one address, which is why the code still needs a barrier.
It also assumes that 32 updates to one address cost 32 times as much as a plain add. This page measures the cost and compares three ways to sum an array at several sizes.
What an atomic promises, and what it does not
The guide is short about it. "Atomic functions perform read-modify-write
operations on shared data, making them appear to execute in a single step",
and atomicAdd "reads a word at a specific address in global or shared
memory, adds a number to it, and writes the result back to the same address"
(https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html
sections 5.4.5 and 5.4.5.1, checked 2026-08-30). One address, one operation,
indivisible, and nothing about other threads.
Two sentences on the same page say what is missing. First: "The atomic
functions described in this section have a memory ordering of
cuda::std::memory_order_relaxed". Second, and this answers the post above:
"Unlike built-in atomic functions, legacy atomic functions only ensure
atomicity and do not introduce synchronization points (fences)."
So an atomic is not a barrier. In that histogram kernel one group of threads
zeroes the shared bins and another adds into them, and atomicity publishes
neither group's stores to the other. That is
__syncthreads()'s job, on
day 14.
An atomic does not fix the order either. Relaxed ordering lets the hardware run the n additions in a different sequence each time. Integer addition still gives the same result.
Float addition is not associative, so a changed order can change the last few bits: "the order in which operations are executed affects the accuracy of the result" (https://docs.nvidia.com/cuda/floating-point/index.html section 3.1, checked 2026-08-30). A sum built from atomics can therefore differ between runs. Day 68 covers floating point determinism.
Why "they take turns" overstates the cost
A published CUDA course says that atomic updates "are serialized at the hardware level", under a diagram captioned "The hardware serializes these operations to update the global accumulator correctly". The same page says it will "Measure performance and discuss the trade-offs", but publishes no measurement (https://github.com/AdepojuJeremy/CUDA-120-DAYS--CHALLENGE/blob/main/daily-updates/day-13-Basic-Atomic-Operations.md , checked 2026-08-30).
It overstates the cost because the compiler gets there first. NVIDIA: "The NVCC compiler now performs warp aggregation for atomics automatically in many cases, so you can get higher performance with no extra effort", where aggregating means "the threads of a warp first compute a total increment among themselves, and then elect a single thread to atomically add the increment to a global counter" (https://developer.nvidia.com/blog/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/ , checked 2026-08-30).
The documented case is a counter. Every lane adds 1 to the same address, so a warp can combine 32 increments with a population count. An array sum is different because each lane holds a different value.
Combining those values needs the warp shuffle
reduction from day 23. Check the machine code to see
whether nvcc inserts one for this float sum. The README gives the
cuobjdump command.
The hardware moved the same way. Maxwell replaced Kepler's
"lock/update/unlock pattern that could be expensive in the case of high
contention" with "native shared memory atomic operations for 32-bit integers"
and native compare-and-swap
(https://docs.nvidia.com/cuda/maxwell-tuning-guide/index.html section
1.4.3.3, checked 2026-08-30). That names integers, not a 32-bit float add, so
whether a shared atomicAdd(float*) on Turing is one instruction or a loop is
something this page measures rather than answers.
Three ways to add up an array
Full program in
code/day26-atomics/atomics.cu. It
sums the same buffer three ways at five sizes, from 4,096 elements to
4,194,304.
The program follows four rules, and the third is the site's own timing rule applied to a kernel that updates its output in place.
Every kernel reads every element once, in the same order. All three walk the input with consecutive threads on consecutive addresses, so the read is coalesced and the bytes pulled out of global memory are identical. The only variable is where the additions land.
The data makes float addition exact, so the check can be ==. Input
values are 0, 1, 2 and 3, so every partial sum is an integer below 2^24 and
exact in float for any addition order. The check uses bitwise equality
against a Kahan double reference, so the table measures
contention, not precision.
A static_assert fails the build if anyone raises the size past that point.
Each timed launch gets its own accumulator. This site's timing rule says an in-place kernel cannot go in a timing loop without a reset between runs. A fresh slot per launch is that reset, and the arithmetic that picks it runs on the host, outside the timed region.
Events, and a warm-up for every kernel timed. Day 9 covers why a host clock around a launch measures the launch.
The naive version is four lines and it is the one the folklore is about:
__global__ void sumGlobalAtomic(const float* __restrict__ in,
float* __restrict__ total, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
atomicAdd(total, in[i]);
}
}
The second gives each block its own accumulator in shared memory and sends one add to the global one. That is privatization, which day 29 applies to histograms, and it is the shape the r/CUDA poster had already written:
__global__ void sumSharedAtomic(const float* __restrict__ in,
float* __restrict__ total, size_t n) {
__shared__ float blockSum;
const unsigned int tid = threadIdx.x;
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
if (tid == 0) {
blockSum = 0.0f;
}
__syncthreads();
if (i < n) {
atomicAdd(&blockSum, in[i]);
}
__syncthreads();
if (tid == 0) {
atomicAdd(total, blockSum);
}
}
The third replaces the shared atomic with the tree reduction from day 24 and keeps the single global add, so it performs the same number of global atomics as the second and none at all inside the block:
__global__ void sumTreeAtomic(const float* __restrict__ in,
float* __restrict__ total, size_t n) {
__shared__ float tile[kThreadsPerBlock];
const unsigned int tid = threadIdx.x;
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
tile[tid] = (i < n) ? in[i] : 0.0f;
__syncthreads();
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]);
}
}
Note. The GB/s column counts the input stream and nothing else, four bytes per element read once. The atomic traffic is left out on purpose: it is the thing the three kernels differ in, so putting it in the denominator would hide the comparison the table exists to make.
Results
Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.
GPU: Tesla T4 (compute capability 7.5)
SMs: 40, threads per block: 256
Sweep: one sum three ways, 10 timed runs after 3 warm-ups
n global ms shared ms tree ms gl GB/s sh GB/s tr GB/s fastest
--------- ---------- ---------- ---------- -------- -------- -------- --------
4096 0.0175 0.0300 0.0047 0.9 0.5 3.5 tree
65536 0.2333 0.0774 0.0097 1.1 3.4 27.0 tree
262144 0.9229 0.2370 0.0259 1.1 4.4 40.5 tree
1048576 3.6827 0.8522 0.0918 1.1 4.9 45.7 tree
4194304 13.1734 1.5241 0.1666 1.3 11.0 100.7 tree
Every row above read the same 4 bytes per element and returned
the same answer bit for bit, checked against a Kahan double
reference with == and not a tolerance. Only the accumulation
strategy changed.
Determinism at n = 4194304, 20 runs each, input that is not
exactly representable. Kahan double reference: 6289253.15
kernel distinct min max max - min
---------------- -------- ---------------- ---------------- ------------
global atomic 20 6289136.5 6289304 167.5
shared atomic 6 6289251.5 6289254 2.5
tree + atomic 3 6289252.5 6289253.5 1
The CUDA 13.0 run preserved the winner and every exact-sum check, but the largest-size timings moved: global, shared and tree changed from 13.1734, 1.5241 and 0.1666 ms to 9.6781, 1.2911 and 0.1424 ms. The tree's advantage there changed from 79x to 68x. The number of distinct floating-point results also varied, as this nondeterminism experiment predicts, while the global atomic still had by far the widest range.
The tree wins at every size, and the gap widens as the input grows. At 4.19 million elements the global-atomic version takes 13.1734 ms and the tree takes 0.1666 ms, which is 79 times faster. Privatising into shared memory first gets most of the way there, 1.5241 ms, without changing the algorithm.
The reason is contention, not atomics being slow. A single global counter serialises every update in the grid; a per-block counter serialises only within a block; a tree serialises almost nothing.
Every strategy returned the same answer bit for bit, checked against a
Kahan double reference with == rather than a tolerance. That is worth stating
because it is only true for this input, where every element is exactly
representable.
Change the input and determinism goes. The run repeats the sweep on values that are not exactly representable and reports how many distinct totals each strategy produced across 20 runs. Float addition is not associative, and an atomic reduction uses the order chosen by the hardware.
The same program can therefore return different sums for the same data. The tree is deterministic because the algorithm fixes its addition order. Day 68 is where that becomes a design constraint.
Run it yourself
A free Colab T4 or any card you own. The build line is in the repo's README:
nvcc -std=c++17 -O3 -arch=sm_75 -o atomics atomics.cu
There is no Compiler Explorer embed on this page. The program allocates a 16 MiB input buffer and makes 270 kernel launches, with 150 inside a timer. That work comes too close to Compiler Explorer's 20 second run limit.
Shared memory is not the cause. The tree kernel's tile is 1 KiB, while a T4 allows 48 KiB per block by default. If you have no GPU, read /setup/learn-cuda-without-a-gpu.
Exercise
Run the sweep as shipped, then change kThreadsPerBlock to 64 and again to
1024, and report each column's ratio to the 256 run at the largest size. Then
say in one sentence why one of the three columns barely moves.
Time: 25 to 40 minutes. Submit: nine ratios, the sentence, and the
fastest column from all three runs.
Check: the harness re-runs the bitwise check at every size and block size, so a kernel that drops elements fails. It prints the size and both values on a mismatch.
The harness checks only the sum. It prints every ratio and the fastest
column whether the sum passes or not. A limit on a ratio would decide the
question that the sweep asks you to test.
Hint 1
Count the atomic operations each kernel sends to the global accumulator, as a
function of n and the block size. Two of the three counts change when you
change the block size. One does not.
Hint 2
At 64 threads per block the grid holds four times as many blocks. What does each block contribute to the global accumulator, and how many barrier steps does a block of 1024 take that 256 does not?
Solution
The naive column barely moves. It performs n global atomics however you cut
the grid up, so the block size changes where those atomics come from and not
how many there are.
The other two get worse at 64. Both send one global atomic per block, so
their count is n / blockSize, and shrinking the block by four multiplies it
by four. Going the other way is not free either: a block of 1024 takes ten
halving steps where 256 takes eight.
What to carry away: the cost of a reduction strategy is the number of contended operations it leaves at each level, and the block size moves work between those levels.
Pitfalls
You use an atomic where you needed a barrier. Atomicity is per address
and per operation: "legacy atomic functions only ensure atomicity and do not
introduce synchronization points (fences)". If one thread has to see what
another wrote, that is __syncthreads() on day 14,
and across blocks it is day 27's fences.
You report a float sum from atomics as reproducible. The second table shows that it is not. Rerun the program and the last bits may change.
A golden-value test with == can pass on your machine and fail in CI for a
reason unrelated to the change under review. Day 68 covers the cost of a
fixed order.
You time an accumulator in a loop without resetting it. Ten timed launches into one address leave you with ten times the answer, and you have stopped checking the kernel too. Give each launch its own accumulator, outside the timed region.
You assume atomics cost 32 times a plain add. The compiler aggregates within a warp in many cases before the hardware sees anything, so the serialized model is an upper bound and often a loose one. Measure it on your card first. Nsight Compute counts the atomic traffic and the README gives the command.
You call it privatized while the block still hammers one address. Moving an accumulator from global to shared memory changes which memory serializes, not whether it does. It pays because a block's contention is cheaper than a grid's and because it cuts the global atomic count by the block size.
You optimise the atomic and the kernel does not move. A sum reads n
floats and writes one scalar, so once the contention is gone its ceiling is
memory bandwidth. Check against day 11's
measured copy number first.
Go deeper
- CUDA Programming Guide 5.4.5, "Atomic Functions", for the operation list and the per-type compute capability rules: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30)
- NVIDIA, "CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics": https://developer.nvidia.com/blog/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/ (checked 2026-08-30)
- NVIDIA, "Floating Point and IEEE 754 Compliance for NVIDIA GPUs", section 3, on why summation order changes the answer: https://docs.nvidia.com/cuda/floating-point/index.html (checked 2026-08-30)
cuda-samples,cpp/2_Concepts_and_Techniques/threadFenceReduction: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/2_Concepts_and_Techniques/threadFenceReduction (checked 2026-08-30)- Programming Massively Parallel Processors, 4th edition, chapter 9, on atomics, privatization and histograms: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0
Next
Day 27 makes compulsory the one thing these kernels never needed: a multi-block reduction that reads its own partial sums is wrong without a fence, and an atomic will not supply the ordering. Day 29 turns the shared accumulator here into a histogram, where the same move has a name, privatization, and where the input rather than the kernel decides what it is worth.