code/day14-syncthreads/syncthreads.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// SPDX-License-Identifier: MIT
//
// Day 14: what __syncthreads() guarantees.
//
// Three kernels over the same data. reverseTileRacy leaves out the barrier
// between the shared write and the shared read. reverseTileSynced is the same
// kernel with that one line put back. reverseTilePrefix returns most of its
// threads before the barrier, which the programming guide's wording allows
// and which does not hang.
//
// Nothing here is timed. The subject is ordering, not speed, so the file has
// no cudaEvent in it and the output carries no number a stopwatch produced.
//
// The racy kernel's mismatch count is printed and never gated. A data race is
// allowed to come out right, and a test demanding that it come out wrong
// would be asserting that undefined behaviour is dependable. The gate is on
// the fixed kernel, which has to be right at every block size on every run.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o syncthreads syncthreads.cu
// Run: ./syncthreads
//
// For the sanitizer pass, rebuild with line numbers and run:
// nvcc -std=c++17 -O3 -lineinfo -arch=sm_75 -o syncthreads syncthreads.cu
// compute-sanitizer --tool racecheck ./syncthreads
//
// Verified 2026-08-30 on a Tesla T4 (compute capability 7.5), driver
// 595.84, CUDA 12.6 (V12.6.85). Transcript: evidence/run-2026-08-30.txt
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda_runtime.h>
// The one error macro. This file is standalone, the way a Compiler Explorer
// embed is, so it carries its own verbatim copy. `err_` carries a trailing
// underscore so it cannot collide with a variable at the call site, and the
// do/while makes the macro one statement so it survives a braceless `if`.
#define CUDA_CHECK(call) \
do { \
cudaError_t err_ = (call); \
if (err_ != cudaSuccess) { \
std::fprintf(stderr, "CUDA error %s:%d: %s: %s\n", __FILE__, \
__LINE__, #call, cudaGetErrorString(err_)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)
// 32 on every GPU this course targets. The built-in `warpSize` is a run-time
// value, so it cannot size an array or appear in a static_assert; this
// constant can do both.
constexpr int kWarpSize = 32;
// The shared tile is sized from the largest block in the sweep, and every
// launch uses only its own blockDim.x entries of it.
//
// The course default block size is 256 and stays the largest here. The four
// sizes are the point of the program: 32 is one warp, and the lesson turns on
// what a single-warp block hides.
constexpr int kMaxThreadsPerBlock = 256;
constexpr int kBlockSizes[] = {32, 64, 128, 256};
constexpr int kBlockSizeCount = 4;
// How many threads per block still take part in reverseTilePrefix. The other
// 160 return above its barrier. A multiple of 32, so whole warps leave and
// the return is warp-uniform.
constexpr unsigned int kPrefixKeep = 96;
// 262,144 plus 611. 611 is 13 x 47, so no block size in the sweep divides the
// element count and every launch ends with a partial tile, which keeps the
// bounds guard live. The size is deliberately modest: compute-sanitizer
// instruments every shared-memory access, so a buffer ten times this one
// turns the racecheck pass in the README into a coffee break.
constexpr size_t kElems = 256ull * 1024ull + 611ull;
// Both sides compute the same permutation of the same small whole numbers, so
// a mismatch here can only be an ordering bug. The tolerance is not zero
// because that is the house rule, not because these inputs need it.
constexpr float kRelTolerance = 1e-5f;
// constexpr so the static_asserts below can call them.
constexpr int blocksFor(size_t n, int threadsPerBlock) {
return static_cast<int>((n + static_cast<size_t>(threadsPerBlock) - 1) /
static_cast<size_t>(threadsPerBlock));
}
constexpr bool everyBlockSizeIsWholeWarps() {
for (int b = 0; b < kBlockSizeCount; ++b) {
if (kBlockSizes[b] % kWarpSize != 0 ||
kBlockSizes[b] > kMaxThreadsPerBlock) {
return false;
}
}
return true;
}
constexpr bool everyBlockSizeLeavesAPartialTile() {
for (int b = 0; b < kBlockSizeCount; ++b) {
if (kElems % static_cast<size_t>(kBlockSizes[b]) == 0) {
return false;
}
}
return true;
}
static_assert(kMaxThreadsPerBlock <= 1024,
"1024 threads per block is the hardware ceiling at every "
"compute capability this course targets");
static_assert(everyBlockSizeIsWholeWarps(),
"every block size must be a whole number of warps and must fit "
"inside the shared tile");
static_assert(everyBlockSizeLeavesAPartialTile(),
"no block size may divide kElems, so the bounds guard runs on "
"every launch instead of never");
static_assert(kBlockSizes[0] == kWarpSize,
"the first row of the sweep has to be a single-warp block; the "
"whole lesson turns on what that row hides");
static_assert(kPrefixKeep % kWarpSize == 0 &&
kPrefixKeep < static_cast<unsigned int>(kMaxThreadsPerBlock),
"reverseTilePrefix must retire whole warps and must retire at "
"least one, or it demonstrates nothing");
static_assert(blocksFor(kElems, kMaxThreadsPerBlock) *
static_cast<size_t>(kMaxThreadsPerBlock) >=
kElems,
"the grid must cover every element");
// Reverses each block-sized tile of `in` into `out`, and gets it wrong.
//
// One thread writes its own element into the tile, then reads the cell its
// mirror wrote. At 256 threads per block thread 0 wants tile[255], which warp
// 7 owns, so almost every read crosses a warp boundary. At 32 threads the
// whole tile belongs to one warp and no read crosses anything.
//
// Memory: one warp's 32 loads from `in` cover 128 contiguous bytes, so the
// global side is coalesced and is not what this kernel is about.
//
// Launch assumption: blockDim.x is a whole number of warps, at most
// kMaxThreadsPerBlock. There is no barrier between the shared write and the
// shared read, which is the bug this lesson exists to show. Do not copy it.
// snippet: racy
__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];
}
}
// end snippet: racy
// The same kernel with the barrier put back. One line, and it is the whole
// difference between the two rows of the table.
//
// One thread writes its own element into the tile, waits until every other
// thread in the block has written theirs, then reads the cell its mirror
// wrote.
//
// Memory: identical to reverseTileRacy. Same loads, same stores, same
// addresses on the global side.
//
// Launch assumption: blockDim.x is a whole number of warps, at most
// kMaxThreadsPerBlock. Every thread in the block reaches the barrier: the
// guard covers the load and the store, never the barrier, and there is no
// return above it.
__global__ void reverseTileSynced(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;
// snippet: barrier
tile[tid] = (i < n) ? in[i] : 0.0f;
__syncthreads();
// end snippet: barrier
const unsigned int src = blockDim.x - 1u - tid;
if (i < n) {
out[i] = tile[src];
}
}
// Reverses only the first `keep` elements of each tile. Threads with
// tid >= keep return before the barrier, and the kernel still finishes.
//
// One thread either retires immediately or writes one element, waits, and
// reads one cell. Whole warps above the keep boundary retire without ever
// reaching the barrier below.
//
// Memory: tile cells at and above `keep` are never written and never read, so
// nothing in here depends on what shared memory held before the block ran.
//
// Launch assumption: keep <= blockDim.x. This is the one kernel in the course
// with a `return` above a __syncthreads(). CUDA-CODE-STYLE.md bans that habit
// and this kernel is why the ban is worth having: a thread that has left is
// not waited for and its writes are not ordered by the barrier, so the only
// safe version is the one where the leavers wrote nothing anybody reads.
// snippet: early-return
__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];
}
}
// end snippet: early-return
// CPU reference for reverseTileRacy and reverseTileSynced.
//
// The block size is a parameter because it is part of the answer here rather
// than a launch detail: reversing inside a 32-wide tile and inside a 256-wide
// tile give different arrays. Written for obvious correctness, not speed. It
// never allocates; the caller owns every buffer.
static void reverseTilesCpu(const float* in, float* out, size_t n,
int blockSize) {
const size_t bs = static_cast<size_t>(blockSize);
for (size_t base = 0; base < n; base += bs) {
for (size_t tid = 0; tid < bs; ++tid) {
const size_t i = base + tid;
if (i >= n) {
break;
}
const size_t j = base + (bs - 1 - tid);
out[i] = (j < n) ? in[j] : 0.0f;
}
}
}
// CPU reference for reverseTilePrefix. Threads at or above `keep` write
// nothing, so those outputs keep the zero the memset put there.
static void reversePrefixCpu(const float* in, float* out, size_t n,
int blockSize, unsigned int keep) {
const size_t bs = static_cast<size_t>(blockSize);
const size_t k = static_cast<size_t>(keep);
for (size_t base = 0; base < n; base += bs) {
for (size_t tid = 0; tid < bs; ++tid) {
const size_t i = base + tid;
if (i >= n) {
break;
}
if (tid >= k) {
out[i] = 0.0f;
continue;
}
const size_t j = base + (k - 1 - tid);
out[i] = (j < n) ? in[j] : 0.0f;
}
}
}
// Counts every element that differs by more than the relative tolerance and
// writes the smallest such index through `firstBad`, or n if they agree
// everywhere. A count and an index rather than a bool: "wrong at 8,191, and
// 4,096 of them are" names the shape of the bug, "wrong" does not.
static size_t countMismatches(const float* got, const float* want, size_t n,
float relTolerance, size_t* firstBad) {
size_t bad = 0;
*firstBad = n;
for (size_t i = 0; i < n; ++i) {
const float scale = (want[i] == 0.0f) ? 1.0f : std::fabs(want[i]);
if (std::fabs(got[i] - want[i]) > relTolerance * scale) {
if (bad == 0) {
*firstBad = i;
}
++bad;
}
}
return bad;
}
// Prints one row of the table.
//
// The last two columns are the diagnosis. For the first wrong element the
// reader is the thread that owns it and the writer is the thread on the other
// side of the tile, and their warp numbers say whether the two are in the
// same warp. A block whose readers and writers always share a warp is a block
// that can hide this bug.
static void printRow(int blockSize, const char* kernel, size_t bad,
size_t firstBad) {
if (bad == 0) {
std::printf("%6d %-18s %12zu %11s %10s %10s\n", blockSize, kernel,
bad, "-", "-", "-");
return;
}
const unsigned int tid =
static_cast<unsigned int>(firstBad % static_cast<size_t>(blockSize));
const unsigned int src = static_cast<unsigned int>(blockSize) - 1u - tid;
std::printf("%6d %-18s %12zu %11zu %10u %10u\n", blockSize, kernel, bad,
firstBad, tid / kWarpSize, src / kWarpSize);
}
int main() {
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
std::printf("GPU: %s (compute capability %d.%d)\n", prop.name, prop.major,
prop.minor);
const size_t bytes = kElems * sizeof(float);
std::printf("n = %zu floats, each block reverses its own tile\n", kElems);
std::printf("%d block sizes, no timing anywhere in this program\n",
kBlockSizeCount);
// Small whole numbers, so every comparison below is exact and a mismatch
// can only be an ordering bug rather than a rounding one.
std::vector<float> h_in(kElems);
std::vector<float> h_out(kElems);
std::vector<float> h_want(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = static_cast<float>(i % 1024);
}
float* d_in = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
// One status variable and one exit. Every path below falls through to the
// frees at the bottom, including the failing ones, because a lesson that
// returns early past its own cudaFree teaches the opposite of day 5.
int status = EXIT_SUCCESS;
std::printf("\n%6s %-18s %12s %11s %10s %10s\n", "block", "kernel",
"mismatches", "first bad", "read warp", "wrote it");
for (int b = 0; b < kBlockSizeCount; ++b) {
const int blockSize = kBlockSizes[b];
const int blocks = blocksFor(kElems, blockSize);
reverseTilesCpu(h_in.data(), h_want.data(), kElems, blockSize);
// The racy kernel first. Its count is reported and never gated: a
// data race may produce the right answer, and a check that demanded
// the wrong one would be asserting that undefined behaviour is
// dependable.
CUDA_CHECK(cudaMemset(d_out, 0, bytes));
reverseTileRacy<<<blocks, blockSize>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
size_t firstBad = kElems;
size_t bad = countMismatches(h_out.data(), h_want.data(), kElems,
kRelTolerance, &firstBad);
printRow(blockSize, "reverseTileRacy", bad, firstBad);
CUDA_CHECK(cudaMemset(d_out, 0, bytes));
reverseTileSynced<<<blocks, blockSize>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
bad = countMismatches(h_out.data(), h_want.data(), kElems,
kRelTolerance, &firstBad);
printRow(blockSize, "reverseTileSynced", bad, firstBad);
// The gate. The fixed kernel has to be right at every block size on
// every run, and this is a real branch rather than an assert because
// CI builds Release, Release defines NDEBUG, and NDEBUG deletes
// assert() out of the build that matters.
if (bad != 0) {
std::fprintf(stderr,
"reverseTileSynced wrong at %zu with %d threads per "
"block: got %.9g, want %.9g, %zu elements differ\n",
firstBad, blockSize, h_out[firstBad], h_want[firstBad],
bad);
status = EXIT_FAILURE;
}
}
// Part 2: a return above the barrier. Every thread at or above
// kPrefixKeep exits before reaching it, and the launch still completes,
// because the barrier waits for the threads that have not exited.
const int prefixBlock = kMaxThreadsPerBlock;
const int prefixBlocks = blocksFor(kElems, prefixBlock);
std::printf(
"\nreverseTilePrefix: %d threads per block, %u reach the barrier\n",
prefixBlock, kPrefixKeep);
reversePrefixCpu(h_in.data(), h_want.data(), kElems, prefixBlock,
kPrefixKeep);
CUDA_CHECK(cudaMemset(d_out, 0, bytes));
reverseTilePrefix<<<prefixBlocks, prefixBlock>>>(d_in, d_out, kElems,
kPrefixKeep);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
size_t prefixFirstBad = kElems;
const size_t prefixBad = countMismatches(
h_out.data(), h_want.data(), kElems, kRelTolerance, &prefixFirstBad);
std::printf(" the launch returned, so the barrier did not deadlock\n");
std::printf(" %zu of %zu elements differ from the prefix reference\n",
prefixBad, kElems);
if (prefixBad != 0) {
std::fprintf(
stderr, "reverseTilePrefix wrong at %zu: got %.9g, want %.9g\n",
prefixFirstBad, h_out[prefixFirstBad], h_want[prefixFirstBad]);
status = EXIT_FAILURE;
}
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
if (status != EXIT_SUCCESS) {
std::fprintf(stderr, "a gated kernel produced a wrong answer\n");
return status;
}
std::printf("\nevery gated kernel matched its reference\n");
return EXIT_SUCCESS;
}