100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryExecution model
CC 7.5

What does __syncthreads() do in CUDA?

A barrier that every thread in a block must reach before any thread passes it, and which also makes shared-memory writes visible to the rest of the block.

Everybody learns the waiting half and stops there, which is where the bugs come from. That half is in the guide plainly: __syncthreads*() "wait until all non-exited threads in the thread block simultaneously reach the same __syncthreads*() intrinsic call in the program or exit". The ordering half is on the same page and is the one a tile kernel is actually buying: the call "strongly happens before" any participating thread leaves the wait. Without that clause your write to a shared memory cell and somebody else's read of it are two events with no order between them, which is the definition of a data race whatever answer comes out. Both promises stop at the thread block, which is also where the shared memory stops. The most-viewed question on this line, 127,990 views, asks whether the scope is the grid, the warp or the block. It is the block, and nothing wider.

The reason people delete it and get away with it is block size. Launch 32 threads and the block is one warp, the store to the tile and the load from it come out of one instruction stream, and the load finds what it wanted. Launch 256 and warp 0 reads cells warp 7 owns while warp 7 is still waiting on its global load, so the order is whatever the scheduler picked that millisecond. Day 14 measured that gap directly, and the shape of the result matters more than the counts: zero wrong at 32 threads, tens of thousands wrong above it. A test at 32 threads a block proves nothing about ordering. Since compute capability 7.0 it does not even prove warp-level lockstep, because independent thread scheduling gave every thread its own program counter and the guide calls warp-synchronous assumptions invalid.

Two follow-up questions come up every time, and both have one-sentence answers. A barrier inside an if is legal when the condition is uniform across the whole block, and undefined otherwise: if (blockIdx.x % 2u == 0u) is fine, if (threadIdx.x < 64u) "may hang or produce unintended side effects". And a thread that returned before the barrier does not deadlock the block, because the guide's wording says non-exited threads "or exit". That is the trap, not the reassurance. The early return quietly drops those threads out of the ordering promise too, so anything they wrote on the way out is no longer ordered against anybody's read below the line. Guard the load and the store, never the barrier.

What it will not do is worth listing, because each item is somebody's bug. It does not synchronize the grid; that needs fences and scopes or a cooperative launch. It does not reach the host; that is cudaDeviceSynchronize(). It does not zero shared memory, so a cell nobody wrote holds the last block's leftovers. And it does not pick a winner when two threads write the same cell, it only decides when the loser becomes visible.

Measured

Day 14 ran two kernels that differ by one line, each block reversing its own tile through shared memory. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, captured 2026-08-30 on the project's verification node. Nothing in that program is timed, because the question is order rather than speed.

Threads per block Mismatches, barrier removed Mismatches, barrier in place
32 0 0
64 17,312 0
128 21,696 0
256 46,656 0

Do not read the three non-zero counts as constants. They record which warp reached the read first, which is scheduling, so another card or another run is free to differ. The result is the zero at 32 and the non-zero above it.

The early-return case ran in the same program. With 256 threads per block and a return above the barrier, only 96 threads reached it, and the launch returned normally rather than hanging. That is the whole hazard in one line: the failure mode of an early return is a wrong answer at full speed, not a deadlock you would notice. Racecheck grades the barrier-less kernel on whether an order exists rather than on which order you happened to get, so it should still object to the 32-thread launch whose output was correct. That has not been run yet; day 14 ships the prediction and the report will land with it. Running it needs a shell, which a hosted web runner does not give you; /setup/learn-cuda-without-a-gpu covers what to do instead.

Diagram

tile-lifecycle, presets no-barrier, block-32-hides-it and early-return, stepping one block through global load, barrier, compute, barrier, store with a per-thread program counter and a scheduling policy you choose.

Alt text: "One block stepping through a tile with the barrier switched off. At 32 threads the write always lands before the read and the answer is right. At 256 threads a lane in warp zero reads a cell that a lane in warp seven has not written yet, and the widget reports the hazard in both cases."

Code

From code/day14-syncthreads/syncthreads.cu. The whole difference between the two rows of that table:

    tile[tid] = (i < n) ? in[i] : 0.0f;
    __syncthreads();                     // delete this line and the table above happens

    const unsigned int src = blockDim.x - 1u - tid;
    if (i < n) {
        out[i] = tile[src];
    }
}

A barrier is a fence between two named accesses. If you cannot say which write it separates from which read, you have added a line rather than fixed a bug: moving it below the read leaves both racing accesses on the same side and changes nothing.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.