Synchronization & Reductions

Why Threads Need to Coordinate

Consider computing the sum of an array. Each thread could add some elements to a total variable—but without coordination, threads overwrite each other's results.

# Naive parallel sum (WRONG - race condition)
@triton.jit
def broken_sum(x_ptr, out_ptr, N):
    pid = tl.program_id(0)
    x = tl.load(x_ptr + pid)
    
    # Every thread reads, adds, writes to same location
    current = tl.load(out_ptr)  # Thread A reads 0
    tl.store(out_ptr, current + x)  # Thread A writes 5
    # Meanwhile Thread B also read 0, writes 3
    # Final result: 3 (not 8!)

The problem is a race condition: multiple threads reading and writing the same memory location. Solutions include:

• Barriers: Force all threads to reach a point before continuing
• Atomics: Hardware-guaranteed read-modify-write operations
• Reduction patterns: Structured algorithms that avoid conflicts

Block-Level Synchronization

__syncthreads() (or tl.debug_barrier() in Triton) is a barrier: all threads in a block must reach it before any can proceed.

# Shared memory pattern requiring barrier
shared = tl.zeros([BLOCK_SIZE], dtype=tl.float32)

# Stage 1: Each thread writes its value
shared[tid] = my_value

# BARRIER: Wait for ALL threads to finish writing
tl.debug_barrier()

# Stage 2: Now safe to read neighbor's values
neighbor = shared[(tid + 1) % BLOCK_SIZE]

# DEADLOCK: Conditional barrier
if tid < 16:
    tl.debug_barrier()  # Only half the threads reach this!
# Threads 16-31 never call barrier → deadlock

# CORRECT: Barrier outside conditional
if tid < 16:
    do_something()
tl.debug_barrier()  # ALL threads reach this

Intra-Warp Communication

Within a warp (32 threads), threads can directly exchange register values using shuffle instructions—no shared memory or barriers needed.

# Warp-level sum reduction using shfl_down
# Each thread starts with its own value
val = my_value

# Log2(32) = 5 steps to reduce 32 values
for offset in [16, 8, 4, 2, 1]:
    val += __shfl_down_sync(0xFFFFFFFF, val, offset)

# After loop: lane 0 has the sum of all 32 values
if lane_id == 0:
    result = val

Parallel Reduction Patterns

A reduction combines N values into 1 using an associative operator (sum, max, min, product). The parallel approach: repeatedly combine pairs until one remains.

Efficient reduction structure:

1. Warp-level: Use shuffles (5 steps for 32 threads, no sync needed)
2. Block-level: Combine warp results in shared memory (need barriers)
3. Grid-level: Either multiple kernel launches or atomics

# Two-level reduction: warp then block
# Step 1: Each warp reduces its 32 values
warp_sum = warp_reduce(my_value)  # Uses shuffles

# Step 2: First thread of each warp writes to shared mem
if lane_id == 0:
    shared[warp_id] = warp_sum

tl.debug_barrier()  # Wait for all warps

# Step 3: First warp reduces the warp sums
if warp_id == 0:
    val = shared[lane_id] if lane_id < num_warps else 0
    block_sum = warp_reduce(val)
    if lane_id == 0:
        output[block_id] = block_sum

Atomic Operations

Atomic operations guarantee that a read-modify-write sequence completes without interference from other threads. The hardware ensures only one thread accesses the location at a time.

# Every thread hits same address
for i in range(1000):
    atomicAdd(total, val[i])
# 1000 serialized atomics!

# Sum locally first
local_sum = 0
for i in range(1000):
    local_sum += val[i]
atomicAdd(total, local_sum)
# Only 1 atomic per thread!

Flexible Synchronization

Cooperative Groups (CUDA 9+) provide flexible thread groupings beyond the fixed block/warp hierarchy. You can synchronize subsets of threads or even across the entire grid.

// Cooperative groups example
#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel() {
    // Get block group (replaces __syncthreads)
    cg::thread_block block = cg::this_thread_block();
    block.sync();
    
    // Get warp-sized tile for shuffle operations
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);
    int sum = cg::reduce(warp, my_val, cg::plus<int>());
    
    // Sub-warp tiles (e.g., 16 threads)
    cg::thread_block_tile<16> half_warp = cg::tiled_partition<16>(block);
}

Hands-On Labs

Further Reading