# Atomics (CUDA)
When many threads update the same location, the read-modify-write has to be indivisible. **Atomics** in CUDA are device functions that perform such an update as one uninterruptible operation, in either [[global-memory-cuda|global]] or [[shared-memory-cuda|shared]] memory.

```c
counter += 1;                 // race: thousands of threads, lost updates
atomicAdd(&counter, 1);       // correct
```

The available operations are `atomicAdd`, `atomicSub`, `atomicExch`, `atomicMin`, `atomicMax`, `atomicInc`, `atomicDec`, `atomicAnd`, `atomicOr`, `atomicXor`, and `atomicCAS`. Each returns the value held *before* the update, which is what makes them usable for building queues and allocators:

```c
int slot = atomicAdd(&head, 1);   // each thread gets a unique slot
buffer[slot] = value;
```

`atomicCAS` is the general primitive the others could be built from, and it is how custom atomic operations get implemented. It is the same compare-and-swap that underpins CPU-side [[atomics]] and lock-free structures.

Contention is the thing to watch. Atomics on a single address from a whole grid serialise, and the throughput collapses. A histogram written naively with one global counter per bin can be slower than the CPU version. The standard fix is a two-level reduction: accumulate into a per-block copy in shared memory, then have one thread per block merge that into global memory.

```c
atomicAdd(&shared_hist[bin], 1);           // fast, contention limited to one block
__syncthreads();
if (threadIdx.x < NBINS)
    atomicAdd(&global_hist[threadIdx.x], shared_hist[threadIdx.x]);
```

This cuts global atomic traffic by roughly the block size. When the operation is a plain sum, a proper tree [[reduction-cuda|reduction]] avoids atomics entirely and is usually faster still.