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 or shared memory.

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:

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.

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 avoids atomics entirely and is usually faster still.