Synchronization & Atomic Operations

Parallel threads that read and write shared data need coordination. Without proper synchronization, you get race conditions — intermittent bugs that produce wrong results silently. This chapter covers every synchronization mechanism CUDA provides, from block-level barriers to lock-free atomic operations.

---

1. The Synchronization Hierarchy

CUDA provides synchronization at three levels, matching the thread hierarchy:

Level	Mechanism	Scope	Cost
Warp	Implicit lockstep (pre-Volta), __syncwarp()	32 threads	~1 cycle
Block	__syncthreads()	All threads in block	~20-50 cycles
Grid	Kernel boundary, atomics, cooperative groups	All blocks	~1000+ cycles

No inter-block synchronization within a kernel (except atomics on global memory). If you need all blocks to synchronize, launch a new kernel.

---

2. Block Synchronization: __syncthreads()

The most commonly used barrier — all threads in a block wait until every thread reaches this point.

When It’s Required

__shared__ float tile[256];

// Phase 1: Write to shared memory
tile[threadIdx.x] = compute(input[idx]);

__syncthreads();  // ← REQUIRED: ensure all writes complete

// Phase 2: Read from shared memory (may read other threads' data)
float val = tile[255 - threadIdx.x];  // safe: all writes visible

Without __syncthreads(), thread 0 might read tile[255] before thread 255 has written it → race condition.

Rules

1. All threads must reach the same __syncthreads() — if some threads skip it (due to conditional), the kernel hangs or produces undefined behavior:

// ❌ WRONG: conditional syncthreads — not all threads reach it
if (threadIdx.x < 128) {
    __syncthreads();  // deadlock: threads 128-255 never arrive
}

// ✅ CORRECT: all threads reach syncthreads, some just skip the work
__syncthreads();
if (threadIdx.x < 128) {
    // do work
}

2. Two syncs for read-after-write patterns — classic tiled GEMM pattern:

// Load tile
tile[threadIdx.x] = data[idx];
__syncthreads();  // sync 1: all loads complete

// Compute with tile
result += tile[...];
__syncthreads();  // sync 2: all computes done before loading next tile

Performance Cost

__syncthreads() costs ~20-50 cycles. In a tiled GEMM with 128 iterations along K, that’s 256 syncs × ~30 cycles = ~7,680 cycles of pure synchronization. This is why minimizing sync frequency matters for performance.

---

3. Warp Synchronization

Pre-Volta: Implicit Lockstep

Before Volta (2017), all threads in a warp executed in perfect lockstep. No explicit sync was needed within a warp — if thread 0 wrote to shared memory, thread 1 could immediately read it.

Volta+: Explicit Warp Sync

Since Volta, threads in a warp can have independent program counters. You must explicitly synchronize:

// Write to shared memory
shared[threadIdx.x % 32] = val;

// ❌ Pre-Volta: this just worked
// ✅ Volta+: explicit sync required
__syncwarp();  // or __syncwarp(mask) for a subset

float neighbor = shared[(threadIdx.x + 1) % 32];

__syncwarp(mask)

Synchronize a subset of threads in a warp:

unsigned mask = 0x0000FFFF;  // only threads 0-15
__syncwarp(mask);  // only these threads must reach this point

Useful when you know only some threads are active (e.g., after a conditional branch).

---

4. Memory Fences

__syncthreads() ensures all threads reach a barrier. Memory fences ensure memory writes are visible to other threads at different scopes.

Three Fence Levels

__threadfence_block();   // writes visible to all threads in same block
__threadfence();         // writes visible to all threads in the device
__threadfence_system();  // writes visible to CPU and other GPUs

When Fences Are Needed

Scenario: A producer block writes data, sets a flag. A consumer block reads the flag, then reads the data.

// Producer block
data[idx] = result;
__threadfence();        // ensure data is visible before flag
flag[blockIdx.x] = 1;  // signal completion

// Consumer block (different block!)
while (flag[producer_id] == 0);  // spin-wait for flag
__threadfence();                  // ensure we see the data
float val = data[idx];            // safe to read

Without __threadfence(), the consumer might see flag=1 but read stale data — the GPU can reorder memory operations.

In practice: Avoid inter-block communication via flags. It’s fragile and hard to debug. Prefer kernel boundaries or atomics.

---

5. Atomic Operations

Atomics perform read-modify-write as a single indivisible operation — no other thread can see an intermediate state.

Common Atomics

atomicAdd(&counter, 1);      // counter += 1
atomicMin(&min_val, val);     // min_val = min(min_val, val)
atomicMax(&max_val, val);     // max_val = max(max_val, val)
atomicExch(&addr, new_val);   // swap addr with new_val
atomicCAS(&addr, expected, desired);  // compare-and-swap

Use Cases

Use case	Atomic	Example
Global counter	atomicAdd	Count elements matching a condition
Global reduction	atomicAdd	Aggregate partial sums from all blocks
Histogram	atomicAdd	Increment bin counts
Lock-free queue	atomicCAS	Producer-consumer across blocks
Global min/max	atomicMin/Max	Find global extremes

Performance

Atomics are slow when contended — many threads updating the same address serialize:

// ❌ High contention: all threads atomic to same address
atomicAdd(&global_sum, local_val);  // 1000+ threads → serialized

// ✅ Two-level reduction: reduce within block first, then one atomic per block
__shared__ float block_sum;
// ... reduce within block using shared memory ...
if (threadIdx.x == 0) {
    atomicAdd(&global_sum, block_sum);  // one atomic per block → low contention
}

Rule: Always reduce contention by doing local (warp/block) reduction first, then a single atomic per block.

Supported Types

Type	atomicAdd	atomicMin/Max	atomicCAS
int	✅	✅	✅
unsigned int	✅	✅	✅
float	✅ (since Volta)	❌	❌
double	✅ (since Volta, slow)	❌	❌
__half	✅ (since Turing)	❌	❌

For unsupported types, use atomicCAS to implement custom atomics:

// Custom atomicMax for float using CAS loop
__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);
}

---

6. Cooperative Groups (Advanced)

Cooperative Groups provide a flexible, type-safe synchronization API that generalizes beyond fixed block/warp boundaries:

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel() {
    // Thread block group
    cg::thread_block block = cg::this_thread_block();
    block.sync();  // same as __syncthreads()

    // Warp-level group
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);
    float sum = cg::reduce(warp, val, cg::plus<float>());  // warp reduce

    // Arbitrary partition
    cg::thread_block_tile<16> half_warp = cg::tiled_partition<16>(block);
    half_warp.sync();
}

Cooperative Groups are the modern way to write synchronization code — they’re more explicit and less error-prone than raw __syncthreads() / __syncwarp().

---

Self-Check

• Explain when __syncthreads() is required and the “all threads must reach” rule
• Explain why two syncs are needed in tiled GEMM (load sync + compute sync)
• Explain the difference between __syncthreads() and __syncwarp()
• Explain when __threadfence() is needed vs __syncthreads()
• Implement a two-level reduction using shared memory + one atomic per block
• Implement a custom atomic using atomicCAS

---

References

• CUDA C++ Programming Guide — Synchronization
• CUDA C++ Programming Guide — Atomic Functions
• Cooperative Groups