GPU Architecture · Cooperation & Optimization

__syncthreads(): the block barrier

How the threads of a block cooperate safely through shared memory — the primitive every tiling kernel is built on.

You've met shared memory (Lesson 5) and the reason it's fast (banks). But cooperation has a catch you already have all the pieces to see. From Lesson 2: the warps of a block interleave in no guaranteed order — the scheduler runs whichever is ready. So when thread A writes a shared slot that thread B reads, who says A ran first? Nobody. That gap is what __syncthreads() closes.

What it is

When a thread calls __syncthreads(), it waits until every thread in its block has reached that same line. Only then does any of them continue. "Nobody moves until everybody's here." It gives two guarantees, and you need both:¹

That second guarantee is the real payload: not just "wait," but "wait — and now everyone can see what everyone wrote."

The race it prevents

Here is the bug, in the tiling shape from Lesson 4. Each thread writes its own shared slot, then reads a neighbour's:

__shared__ float tile[256];
tile[threadIdx.x] = input[gid];      // each thread writes ITS slot
// ✗ RACE if you read here — the neighbour's warp may not have run yet
float x = tile[threadIdx.x + 1];     // reads a NEIGHBOUR's slot

Thread 0 reads tile[1] — but the warp owning tile[1] might not have been scheduled yet, so the slot is uninitialised garbage. Wrong answers that change with occupancy. The barrier turns "warps run in unknown order" into clean phases:

The classic pattern needs a barrier on both sides

A real tiling loop usually has two syncs per iteration, each guarding an opposite hazard:

tile[threadIdx.x] = input[...];                 // LOAD phase
__syncthreads();                                // ① don't READ until the tile is fully written
sum += tile[threadIdx.x] + tile[threadIdx.x ^ 1]; // COMPUTE phase (reads tile)
__syncthreads();                                // ② don't OVERWRITE until everyone finished reading
tile[threadIdx.x] = input[...next...];          // next LOAD would clobber otherwise

• ① read-after-write: consumers must not start until producers finish.
• ② write-after-read: producers must not overwrite until consumers finish. Drop this and a fast warp clobbers the tile while a slow warp is still reading it.

Why it's block-only — the question you parked

There is deliberately no cross-block __syncthreads(). Recall Lesson 3: blocks are independent — the GigaThread Engine runs them in any order, in waves, on different SMs; two blocks may not even be resident at the same time. Asking block A to wait for block B, when B hasn't launched yet, would deadlock by construction. So the only grid-wide barrier is the end of the kernel: if phase 2 needs all blocks of phase 1 done, you split it into two kernel launches.

The footgun: all-or-none

Because the barrier waits for every thread, all threads must reach the same __syncthreads() — or none. Hide it behind a thread-divergent branch and you break it:

if (threadIdx.x < 100) {
    __syncthreads();   // ✗ threads 100+ never arrive → barrier never completes → hang
}

The threads that skip it never increment the counter, so the waiting ones hang forever (pre-Volta: hard deadlock; generally: undefined behaviour). Rule: a barrier inside a conditional is safe only if the condition is uniform across the whole block (e.g. based on blockIdx) — never on a threadIdx-divergent path. compute-sanitizer --tool synccheck catches violations.²

Run it yourself

The companion program reverses each block's segment through shared memory — the reversed write reads a neighbour's slot, so it needs the barrier. Build it both ways:

# correct — barrier present
nvcc -arch=sm_120 02-syncthreads-reverse.cu -o rev && ./rev
# barrier removed — a data race (may pass by luck, or FAIL)
nvcc -arch=sm_120 -DSKIP_SYNC 02-syncthreads-reverse.cu -o rev_bad && ./rev_bad
# let the tool find the race the naked eye can't:
compute-sanitizer --tool racecheck ./rev_bad

./rev prints PASS. The rev_bad build may still pass on a quiet GPU — races are nondeterministic — but racecheck will report the shared-memory hazard regardless. That's the lesson: correctness you can't see is exactly why the barrier and the sanitizer both exist. File: code/02-syncthreads-reverse.cu.

Check yourself

From memory — one click locks each answer.

Notes & citations

1. CUDA C++ Programming Guide, "Synchronization Functions" — __syncthreads() waits until all threads in the block reach it and acts as a memory fence for shared and global accesses made by the block. docs.nvidia.com
2. Same source: __syncthreads() is permitted in conditional code only if the condition evaluates identically across the entire thread block; otherwise execution is likely to hang or behave unexpectedly.