gpu atomic operation
**GPU Atomic Operations** are the **hardware-supported read-modify-write instructions that guarantee indivisible updates to shared memory locations even when thousands of GPU threads access the same address simultaneously** — essential for reductions, histograms, counters, and lock-free data structures on GPUs, where the massive thread parallelism makes unprotected concurrent writes catastrophically incorrect, but where naive use of atomics creates severe contention bottlenecks that can reduce GPU throughput by 10-100×.
**Why Atomics on GPU**
- 10,000+ concurrent threads → many threads may write to same memory location.
- Without atomic: Thread A reads value 5, Thread B reads 5, both write 6 → should be 7 (lost update).
- With atomic: atomicAdd(&counter, 1) → hardware serializes → correct result guaranteed.
- GPU hardware: Dedicated atomic units in L2 cache and shared memory.
**Available Atomic Operations (CUDA)**
| Operation | Function | Supported Types |
|-----------|----------|----------------|
| Add | atomicAdd(addr, val) | int, float, double (sm_60+) |
| Subtract | atomicSub(addr, val) | int |
| Min/Max | atomicMin/atomicMax | int, unsigned int |
| Exchange | atomicExch(addr, val) | int, float |
| Compare-and-swap | atomicCAS(addr, compare, val) | int, unsigned long long |
| Bitwise | atomicAnd/Or/Xor | int, unsigned int |
| Increment | atomicInc(addr, val) | unsigned int |
**Performance Characteristics**
```cuda
// Worst case: All threads atomic to same address
atomicAdd(&global_sum, local_val); // 10000 threads → serialized → very slow
// Better: Warp-level reduction first, then one atomic per warp
float warp_sum = warpReduceSum(local_val); // 32 threads → 1 value
if (lane_id == 0)
atomicAdd(&global_sum, warp_sum); // 32× fewer atomics
// Best: Block-level reduction, then one atomic per block
float block_sum = blockReduceSum(local_val); // 256 threads → 1 value
if (threadIdx.x == 0)
atomicAdd(&global_sum, block_sum); // 256× fewer atomics
```
**Contention Impact**
| Pattern | Threads per address | Throughput |
|---------|-------------------|------------|
| No contention (unique addresses) | 1 | ~500 Gops/s |
| Low contention (per-warp) | 32 | ~50 Gops/s |
| Medium contention (per-block) | 256 | ~10 Gops/s |
| High contention (all same) | 10000+ | ~0.1 Gops/s |
**Shared Memory vs. Global Memory Atomics**
- Shared memory atomics: ~5 ns (same SM, fast path).
- Global memory atomics: ~50-200 ns (L2 cache, may serialize across SMs).
- Strategy: Do atomics in shared memory → final result atomic to global.
**Histogram Example**
```cuda
__global__ void histogram(int *data, int *hist, int n) {
__shared__ int local_hist[256]; // Local histogram per block
if (threadIdx.x < 256) local_hist[threadIdx.x] = 0;
__syncthreads();
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n)
atomicAdd(&local_hist[data[idx]], 1); // Shared mem atomic (fast)
__syncthreads();
// Merge to global histogram
if (threadIdx.x < 256)
atomicAdd(&hist[threadIdx.x], local_hist[threadIdx.x]); // One atomic per bin per block
}
```
**CAS-Based Custom Atomics**
```cuda
// Custom atomicMax for float (not natively supported on all archs)
__device__ float atomicMaxFloat(float *addr, float val) {
int *addr_as_int = (int*)addr;
int old = *addr_as_int, assumed;
do {
assumed = old;
old = atomicCAS(addr_as_int, assumed,
__float_as_int(fmaxf(val, __int_as_float(assumed))));
} while (assumed != old);
return __int_as_float(old);
}
```
GPU atomic operations are **the correctness foundation for concurrent GPU data structures** — while their naive use creates devastating serialization bottlenecks that negate GPU parallelism, the hierarchical reduction pattern (warp → block → global) transforms atomics from a performance liability into a practical tool that enables histograms, counters, and dynamic data structures to work correctly at GPU scale with acceptable overhead.