barrier synchronization
**Barrier Synchronization** is the **fundamental parallel coordination primitive where all participating threads or processes must arrive at a designated point before any can proceed past it** — ensuring a consistent global state at synchronization points, implemented through algorithms ranging from simple centralized counters to sophisticated tree-based and butterfly barriers that scale to thousands of threads while minimizing contention and latency.
**Why Barriers**
- Parallel phases: Phase 1 (compute) → barrier → Phase 2 (exchange) → barrier → Phase 3 (compute).
- Without barrier: Thread A starts phase 2 while thread B is still in phase 1 → reads stale data.
- Barrier guarantees: All threads completed phase 1 before any enters phase 2.
- Common uses: Iterative solvers, BSP model, GPU __syncthreads(), MPI_Barrier().
**Centralized Barrier (Simple Counter)**
```c
// Simplest barrier: atomic counter + spinning
typedef struct {
atomic_int count;
atomic_int sense;
int num_threads;
} barrier_t;
void barrier_wait(barrier_t *b, int *local_sense) {
*local_sense = !(*local_sense); // Flip local sense
if (atomic_fetch_add(&b->count, 1) == b->num_threads - 1) {
// Last thread to arrive → release all
atomic_store(&b->count, 0);
atomic_store(&b->sense, *local_sense);
} else {
// Spin until sense flips
while (atomic_load(&b->sense) != *local_sense) { }
}
}
```
- Problem: All threads contend on single counter → O(P) serialization.
- All threads spin on same variable → cache line bouncing on multi-socket systems.
**Tree Barrier (Logarithmic)**
```
[Root]
/ \
[N0] [N1]
/ \ / \
[T0] [T1] [T2] [T3]
Arrival (up): T0→N0, T1→N0, T2→N1, T3→N1 → N0→Root, N1→Root
Release (down): Root→N0,N1 → N0→T0,T1 → N1→T2,T3
```
- O(log P) steps instead of O(P).
- Each node only communicates with parent/children → reduced contention.
- Natural fit for NUMA: Tree structure matches socket/core topology.
**Butterfly (Tournament) Barrier**
```
Step 0: T0↔T1, T2↔T3 (pairs at distance 1)
Step 1: T0↔T2, T1↔T3 (pairs at distance 2)
After log₂(P) steps: All threads know everyone has arrived.
```
- O(log P) steps, all threads active every step → maximum parallelism.
- No single bottleneck node → better than tree for large P.
- Each step: Thread i synchronizes with thread i XOR 2^step.
**GPU Barriers**
| Scope | Mechanism | Latency |
|-------|-----------|--------|
| Warp (32 threads) | __syncwarp() | ~1 cycle (implicit in SIMT) |
| Thread block (up to 1024) | __syncthreads() | ~20-40 cycles |
| Grid (all blocks) | cooperative_groups::grid_group::sync() | ~1000+ cycles |
| Multi-GPU | NCCL barrier / cudaDeviceSynchronize | ~µs |
**Barrier Performance on Multi-Socket Servers**
| Algorithm | 2 threads | 64 threads | 256 threads |
|-----------|----------|-----------|------------|
| Centralized | 50 ns | 2 µs | 15 µs |
| Tree (degree-2) | 50 ns | 400 ns | 800 ns |
| Butterfly | 50 ns | 300 ns | 600 ns |
| MCS (scalable) | 50 ns | 350 ns | 650 ns |
**Sense-Reversing Barrier**
- Problem: Reusing barrier immediately after release → threads from previous barrier mix with next.
- Solution: Each barrier invocation uses opposite sense (true/false) → threads only wake on matching sense.
- Eliminates need to reset barrier state between consecutive uses.
Barrier synchronization is **the heartbeat of bulk-synchronous parallel computing** — every iterative solver, every GPU kernel with shared memory cooperation, and every MPI collective operation depends on efficient barriers to enforce ordering between computation phases, making barrier algorithm choice a critical performance factor for any parallel application that synchronizes more than a few dozen threads.