Synchronization & Atomics

⚠️ __syncthreads() inside divergent code (where some threads of the block don't reach it) is undefined behavior (deadlock/corruption).

12.1 Synchronization levels

all threads of a block (barrier + memory visibility within the block) threads of a warp barrier with a reduction of a predicate host waits for all device work cudaStreamSynchronize(s) / events host waits for a stream / event implicit global sync between kernels in the same stream cooperative groups grid.sync() whole grid (cooperative launch, §23)

__threadfence(), __threadfence_block(), __threadfence_system() memory ordering (not barriers)

⚠️ __syncthreads() inside divergent code (where some threads of the block don't reach it) is undefined behavior (deadlock/corruption).

12.2 Race conditions

__global__ void racy(int* counter) {
    (*counter)++;            // WRONG: read-modify-write races between threads
}
__global__ void safe(int* counter) {
    atomicAdd(counter, 1);   // correct, but serializes on one address (contention)
}

12.3 Atomic functions

cuda::atomic_ref<T, cuda::thread_scope_device> (libcu++) gives C++-style atomics with explicit memory order.

atomicAdd(&x, v);        // int, unsigned, float, double (CC 6.0+), half2, bfloat16
atomicSub, atomicExch, atomicMin, atomicMax, atomicInc, atomicDec
atomicAnd, atomicOr, atomicXor
atomicCAS(&x, compare, val)     // compare-and-swap: build any atomic operation
atomicAdd_block(...), atomicAdd_system(...)   // scope variants

// Custom atomic via CAS: atomic max for float
__device__ float atomicMaxFloat(float* addr, float val) {
    int* ai = (int*)addr;
    int old = *ai, assumed;
    do {
        assumed = old;
        if (__int_as_float(assumed) >= val) break;
        old = atomicCAS(ai, assumed, __float_as_int(val));
    } while (assumed != old);
    return __int_as_float(old);
}

12.4 Reducing contention

Use per-block partial results in shared memory (privatization), then one atomic per block; warp-aggregated atomics (one atomic per warp after a warp reduction).