100 DAYS OF GPU Course preview
Day 14Module 2
in-technical-review

What __syncthreads() guarantees

Someone working through GPU Puzzles, on a tile kernel with a barrier left out:

"I notice that it passes regardless if I do a syncthreads between reading the shared memory and modifying them for the next iteration. Why the absence of that sync does not cause any issue?"

https://github.com/srush/GPU-Puzzles/issues/14 (checked 2026-08-29)

Their test reported the right output, but the program still had a data race. A 32-thread block has one warp, so the missing barrier may not affect that run. The same kernel fails with 256 threads.

This page says what the barrier promises, shows a tile kernel that is right at 32 threads a block and wrong at 256, and proves the race with a tool that does not care which answer you got.

What the barrier promises, and to whom

__syncthreads() makes two promises, and people usually only know the first one.

It waits. The programming guide says: "__syncthreads*() wait until all non-exited threads in the thread block simultaneously reach the same __syncthreads*() intrinsic call in the program or exit." (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html , checked 2026-08-30.) The scope is the thread block, not the grid or warp.

The question "Does __syncthreads() synchronize all threads in the grid or just the threads in the current warp or block?" had 127,990 views (https://stackoverflow.com/questions/15240432/does-syncthreads-synchronize-all-threads-in-the-grid , checked 2026-08-29). The answer is one block.

It orders memory. The same page says: "__syncthreads*() provide memory ordering among participating threads: the call to __syncthreads*() intrinsics strongly happens before ... any participating thread is unblocked from the wait or exits." A tile kernel needs this order between one thread's write to shared memory and another thread's read of the same cell. Without that order, the accesses form a data race.

Both guarantees stop at the block, as does shared memory. The barrier says nothing about another block or about the initial contents of shared memory. It also does not make two writes to one cell safe.

Why a 32-thread block hides it

A 32-thread block has one warp. Its 32 lanes issue the tile store before the tile load in this run, so the load gets the expected value.

A 256-thread block has eight warps. Warp 0 can read cells owned by warp 7 while warp 7 still waits on its global load. The scheduler does not guarantee an order between those warps.

The 32-thread launch can produce the right output without making the code safe. Since compute capability 7.0, independent thread scheduling means that code cannot assume lockstep within a warp: "Warp-synchronous code assumes that threads in the same warp execute in lockstep at every instruction, but the ability for threads to diverge and reconverge at sub-warp granularity makes such assumptions invalid" (https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-kernel-programming.html , checked 2026-08-29).

The widget's block-32-hides-it preset shows this case: it produces the right answer but still marks the access as undefined. A test at 32 threads per block does not prove correct ordering.

Use __syncwarp() for warp scope and __syncthreads() for block scope.

The barrier in an if, and the thread that already left

Two questions cluster here, and both have a one-sentence answer in the guide.

A barrier inside an if is legal when the condition is the same for every thread in the block. "The __syncthreads*() intrinsics are permitted in conditional code, but only if the condition evaluates uniformly across the entire thread block", and otherwise "execution may hang or produce unintended side effects" (same page as above, checked 2026-08-30). The rule is about the condition, not the if.

if (blockIdx.x % 2u == 0u) {  // same for every thread in the block: legal
    __syncthreads();
}

if (threadIdx.x < 64u) {      // different per thread: undefined
    __syncthreads();
}

A thread that returned before the barrier does not deadlock the block. This is the other half of the wording above: the barrier waits for the non-exited threads, "or exit". A learner asking about it put the confusion precisely: "The documentation states that __syncthreads() must be called by every thread in the block or else it will lead to a deadlock, but in practice I have never experienced such behavior." (https://stackoverflow.com/questions/6666382/can-i-use-syncthreads-after-having-dropped-threads , checked 2026-08-29.)

It does not deadlock. An early return removes those threads from both guarantees, so their earlier writes are not ordered against reads below the barrier.

This is why day 5 used if (i < n) { ... } instead of if (i >= n) return;.

One line apart

The full program is in code/day14-syncthreads/syncthreads.cu. Each block reverses one tile through shared memory and needs one barrier. This isolates the load, barrier, and read sequence from day 13.

Three rules hold it together.

Two kernels that differ by one line. Same input, same launch configuration, same global access pattern. The only variable in the table is the barrier, so a difference in the answer has one candidate cause.

The broken kernel is reported and never gated. A data race is allowed to produce the right answer. A check that demanded the wrong one would be asserting that undefined behaviour is dependable. The gate is on the fixed kernel, which has to match the CPU reference at all four block sizes on every run.

Nothing is timed. No cudaEvent, no stopwatch, no number in the output that a clock produced. This lesson is about order, and day 42 is where a barrier's cost turns up as a stall reason.

Here is the broken one:

__global__ void reverseTileRacy(const float* __restrict__ in,
                                float* __restrict__ out, size_t n) {
    __shared__ float tile[kMaxThreadsPerBlock];

    const unsigned int tid = threadIdx.x;
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;

    tile[tid] = (i < n) ? in[i] : 0.0f;

    // The missing __syncthreads() belongs on this line.

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

The fix is the line the comment is sitting on:

    tile[tid] = (i < n) ? in[i] : 0.0f;
    __syncthreads();

And the early return, which the program runs because it is defined and finishes:

__global__ void reverseTilePrefix(const float* __restrict__ in,
                                  float* __restrict__ out, size_t n,
                                  unsigned int keep) {
    __shared__ float tile[kMaxThreadsPerBlock];

    const unsigned int tid = threadIdx.x;
    if (tid >= keep) {
        return;
    }

    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
    tile[tid] = (i < n) ? in[i] : 0.0f;
    __syncthreads();

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

This is the only course kernel with a return above a barrier. The threads that return do not write any data that another thread reads. If they did, the pattern would be unsafe.

Results

Measured. 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; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5)
n = 262755 floats, each block reverses its own tile
4 block sizes, no timing anywhere in this program

 block  kernel               mismatches    first bad   read warp    wrote it
    32  reverseTileRacy               0            -           -           -
    32  reverseTileSynced             0            -           -           -
    64  reverseTileRacy           17312          288           1           0
    64  reverseTileSynced             0            -           -           -
   128  reverseTileRacy           21696            0           0           3
   128  reverseTileSynced             0            -           -           -
   256  reverseTileRacy           46656          128           4           3
   256  reverseTileSynced             0            -           -           -

reverseTilePrefix: 256 threads per block, 96 reach the barrier
  the launch returned, so the barrier did not deadlock
  0 of 262755 elements differ from the prefix reference

every gated kernel matched its reference

The race is invisible at 32 threads and unmistakable above it. With a 32-thread block the unsynchronised kernel produced zero mismatches. At 64 it produced 17,312. At 256, 46,656.

A 32-thread block is one warp, so this run issues the write before the read. The missing barrier does not change its output. A test with this block size can therefore pass even though a larger block produces wrong results.

The mismatch counts depend on warp scheduling, so they can change between GPUs and runs. The key result is zero mismatches at 32 threads and non-zero counts at larger sizes.

The README asks you to run the program several times. The CUDA 13.0 re-run produced 17,632, 25,984, and 48,928 mismatches at 64, 128, and 256 threads. Those different counts kept the same pass and fail pattern.

The early-return case does not deadlock. In reverseTilePrefix, only 96 of 256 threads reach the barrier, and the launch returns normally. The barrier waits only for threads that have not exited.

The barrier does not order writes made by the exited threads against later reads. This can cause wrong output without a hang.

Run it yourself

A free Colab T4 or any card you own. The program is small enough for Compiler Explorer, but the exercise is not: compute-sanitizer needs a shell, which a shared web runner does not give you. Build line from the repo's README, with line numbers on so the sanitizer can name a source line:

nvcc -std=c++17 -O3 -lineinfo -arch=sm_75 -o syncthreads syncthreads.cu

Hardware. Whether compute-sanitizer runs on a free hosted tier is an open question in this project, not a settled one. Its manual documents a permission requirement only for Jetson and Drive Tegra devices, which is a different restriction from the ERR_NVGPUCTRPERM counter block that stops Nsight Compute. Until each tier is checked, read the report shipped beside this lesson, or start at /setup/learn-cuda-without-a-gpu.

Exercise

Delete the __syncthreads() from reverseTileSynced and run three times. Put one back below the read instead of above it, and run again. Then put it where it belongs and run compute-sanitizer --tool racecheck --racecheck-report all on the middle version and on the fixed one, so Compute Sanitizer grades a kernel the program passed.

Time: 25 to 40 minutes. Submit: which block sizes failed in step 1, whether step 2 changed anything, and one sentence saying why racecheck objects to a 32-thread launch whose output was right.

Check: two gates, both real branches returning EXIT_FAILURE rather than an assert, because CI builds Release and NDEBUG deletes an assert out of the build that matters. A failure names the kernel, the block size, the first wrong index, got, wanted and the count.

Green at 32 and red at 256 is the signature of a missing barrier. Harness contract at /reference/harness.

Hint 1

A barrier is not something a kernel has. It is something that sits between two accesses. Name the two in this kernel that need an order between them, then say where the line goes relative to both.

Hint 2

At 32 threads a block, how many warps are in the block, and which warp writes the tile cell that thread 0 reads? Ask the same two questions at 256 threads and the answer changes.

Solution

The barrier must sit between the tile write and tile read shown above. A barrier elsewhere does not order those two accesses.

The step 2 kernel contains __syncthreads() but remains wrong at the same block sizes. Both racing accesses are on the same side of the barrier. Racecheck reports the same hazard as the version with no barrier.

At 32 threads a block the tile belongs to one warp, so both accesses come out of one instruction stream and the read finds what it wanted. racecheck objects anyway: it checks whether an order exists, not which order you got.

A barrier orders specific accesses. Name the write and read that need the order, then place the barrier between them.

Pitfalls

Your kernel passes at 32 threads a block and you ship it. One warp has nothing to interleave with, so a missing barrier costs nothing and proves nothing. Test at the block size you actually launch, and run racecheck at every size. Day 62 makes that a habit.

You add a __syncthreads() and the answer is still wrong. The barrier orders the accesses on either side of it, not the kernel as a whole. Find the write and the read that race, and put the line between them.

You forget the second barrier in a loop. The GPU Puzzles question at the top of this page is this one: the next iteration's write to a tile cell races with this iteration's read of it. A reloaded tile needs a barrier after the load and another after the last read. Day 16 shows the wrong result when a matrix multiply omits the second barrier.

You return early above a barrier. It will not hang. It removes those threads from the barrier's guarantee, so whatever they wrote before leaving is no longer ordered against anybody's read. Guard the loads and the stores, never the barrier.

You put the barrier inside a thread-dependent if. The condition has to be uniform across the block or "execution may hang or produce unintended side effects". if (blockIdx.x % 2u == 0u) is fine; if (threadIdx.x < 64u) is not. Day 65 diagnoses the hang.

You expect the barrier to synchronize the grid. It is block scope, like the shared memory it exists to order. Ordering between blocks needs fences and scopes (day 27) or a cooperative launch (day 28).

Go deeper

Next

Day 15 keeps the tile and changes only the addresses inside it, which costs bandwidth rather than correctness. Day 16 puts the tile in a loop, where one barrier stops being enough. Both sit in module 2, and day 62 uses racecheck on a less obvious race.