CUDA Cheatsheet

Synchronization

Use this CUDA reference while you build software engineering projects, review code for technical interview prep, or polish examples for a software engineer resume.

Thread-Block Barrier: `__syncthreads`

Synchronizes all threads in a block. No thread may pass until every thread has reached the barrier.

__global__ void reverseBlock(float* data, int n) {
    __shared__ float tile[256];
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // Load phase
    if (i < n) tile[threadIdx.x] = data[i];
    __syncthreads();   // ensure all loads are done before reads

    // Read opposite end
    if (i < n) data[i] = tile[blockDim.x - 1 - threadIdx.x];
}

Gotcha: All threads in the block must reach __syncthreads(). Placing it inside a conditional that not all threads enter causes deadlock or undefined behavior.

// WRONG — only some threads hit syncthreads
if (threadIdx.x < 128) {
    tile[threadIdx.x] = val;
    __syncthreads();    // threads 128-255 never arrive → deadlock
}

// CORRECT — all threads participate
tile[threadIdx.x] = (threadIdx.x < 128) ? val : 0.0f;
__syncthreads();

Variants of `__syncthreads`

Function	Description
`__syncthreads()`	Full block barrier (memory + execution)
`__syncthreads_count(pred)`	Returns number of threads where pred ≠ 0
`__syncthreads_and(pred)`	Returns non-zero if pred ≠ 0 for ALL threads
`__syncthreads_or(pred)`	Returns non-zero if pred ≠ 0 for ANY thread

int n_active = __syncthreads_count(my_condition);   // reduction + barrier

Warp-Level Synchronization

Warps are naturally synchronized (SIMT), but independent thread scheduling (CC ≥ 7.0 / Volta) means you cannot assume warp-level convergence without an explicit barrier.

// CC ≥ 7.0: use __syncwarp() for warp-level barrier
__global__ void warpReduce(float* data) {
    __shared__ float smem[32];
    int lane = threadIdx.x & 31;

    smem[lane] = data[threadIdx.x];
    __syncwarp();           // all 32 lanes visible to each other

    // Warp reduction
    for (int s = 16; s >= 1; s >>= 1) {
        if (lane < s) smem[lane] += smem[lane + s];
        __syncwarp();       // barrier after each step
    }
}

// __syncwarp with mask (only sync the lanes in the mask)
unsigned mask = 0xFFFFFFFF;
__syncwarp(mask);

Atomic Operations

Atomic operations read-modify-write globally without race conditions. All return the old value.

Integer atomics

Function	Operation
`atomicAdd(ptr, val)`	`*ptr += val`
`atomicSub(ptr, val)`	`*ptr -= val`
`atomicExch(ptr, val)`	`*ptr = val`
`atomicMin(ptr, val)`	`ptr = min(ptr, val)`
`atomicMax(ptr, val)`	`ptr = max(ptr, val)`
`atomicInc(ptr, val)`	`ptr = (ptr >= val) ? 0 : *ptr+1`
`atomicDec(ptr, val)`	`ptr = (ptr == 0		ptr > val) ? val : ptr-1`
`atomicCAS(ptr, compare, val)`	`ptr = (ptr==compare) ? val : *ptr`
`atomicAnd(ptr, val)`	`*ptr &= val`
`atomicOr(ptr, val)`	`*ptr	= val`
`atomicXor(ptr, val)`	`*ptr ^= val`

Supported address spaces: __global__, __shared__, __managed__.

__global__ void histogram(int* hist, const int* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) atomicAdd(&hist[data[i]], 1);
}

Float atomics

// atomicAdd(float*) — supported on CC ≥ 2.0
atomicAdd(&d_sum, myFloat);

// atomicAdd(double*) — CC ≥ 6.0
atomicAdd(&d_dsum, myDouble);

// atomicAdd(__half2*) — CC ≥ 7.0 (two 16-bit in one operation)
atomicAdd((__half2*)ptr, __halves2half2(a, b));

CAS-based custom atomics

// atomicMax for float (no native float max atomic pre-CC 8.x)
__device__ float atomicMaxFloat(float* addr, float val) {
    int* iaddr = (int*)addr;
    int  old   = *iaddr, assumed;
    do {
        assumed = old;
        old = atomicCAS(iaddr, assumed,
                        __float_as_int(fmaxf(val, __int_as_float(assumed))));
    } while (assumed != old);
    return __int_as_float(old);
}

Memory Fences

Control the visibility of memory operations to other threads; they are NOT execution barriers.

Fence	Scope	Effect
`__threadfence_block()`	Block	All prior stores visible to block threads
`__threadfence()`	Device	All prior stores visible to all device threads
`__threadfence_system()`	System	Visible to host and all devices (for mapped/peer memory)

__global__ void producer(int* flag, float* data) {
    data[0] = 42.0f;
    __threadfence();       // ensure data write is visible before flag
    atomicExch(flag, 1);   // signal consumer
}

Cooperative Groups

The modern, composable synchronization API (CUDA 9+, all supported architectures). Groups represent sets of cooperating threads at any granularity. Grid-wide sync (cg::this_grid().sync()) additionally requires a cooperative launch and CC ≥ 6.0.

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

__global__ void cgDemo(float* data) {
    // Pre-defined groups
    auto block = cg::this_thread_block();   // all threads in block
    auto warp  = cg::tiled_partition<32>(block); // one warp
    auto tile4 = cg::tiled_partition<4>(block);  // 4-thread tile

    // Barrier
    block.sync();
    warp.sync();
    tile4.sync();

    // Collective operations on a tile
    float val = data[block.thread_rank()];
    float sum = cg::reduce(warp, val, cg::plus<float>());
    float mx  = cg::reduce(tile4, val, cg::greater<float>());
}

Cooperative group types

Group	How obtained	Sync scope
`thread_block`	`cg::this_thread_block()`	Block
`grid_group`	`cg::this_grid()`	Grid (cooperative launch)
`multi_grid_group`	`cg::this_multi_grid()`	Multi-device (cooperative)
`coalesced_group`	`cg::coalesced_threads()`	Active threads in warp
`thread_block_tile<N>`	`cg::tiled_partition<N>(parent)`	N threads (N: 1,2,4,8,16,32)
`cluster_group`	`cg::this_cluster()`	Cluster (CC ≥ 9.0)

Cooperative group collective operations

Operation	Description
`cg::reduce(g, val, op)`	Reduce across group
`cg::inclusive_scan(g, val, op)`	Inclusive prefix scan
`cg::exclusive_scan(g, val, op)`	Exclusive prefix scan
`g.shfl(val, src)`	Shuffle — broadcast from src
`g.shfl_up(val, delta)`	Get `val` from lane `rank − delta` (lower lane)
`g.shfl_down(val, delta)`	Get `val` from lane `rank + delta` (higher lane)
`g.any(pred)`	True if any lane's pred
`g.all(pred)`	True if all lanes' pred
`g.ballot(pred)`	Bitmask of which lanes satisfy pred
`g.match_any(val)`	Bitmask of lanes with matching val
`g.match_all(val, &pred)`	All lanes same value?

Host-Side Device Synchronization

cudaDeviceSynchronize();              // block CPU until ALL GPU work is done
cudaStreamSynchronize(stream);        // block CPU until stream is done
cudaEventSynchronize(event);          // block CPU until event is recorded

// Non-blocking query
cudaError_t err = cudaStreamQuery(stream);
if (err == cudaSuccess)        { /* done */ }
else if (err == cudaErrorNotReady) { /* still running */ }

Warp Vote Functions (CC ≥ 7.0: use `_sync` variants)

unsigned mask = __activemask();
int any_true = __any_sync(mask, predicate);   // 1 if any lane satisfies pred
int all_true = __all_sync(mask, predicate);   // 1 if all lanes satisfy pred
unsigned ballot = __ballot_sync(mask, predicate);  // bitmask of satisfying lanes

Deprecated (CC < 7.0): __any(pred), __all(pred), __ballot(pred) — replace with _sync variants everywhere.