code/day62-racecheck/racecheck.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 62: finding races with racecheck and synccheck.
//
// Two bugs, planted on purpose, each carrying the fixed version beside it:
//
// scanTileRacy day 31's in-place tile scan with the second
// barrier missing. Racecheck's target.
// scanTileFixed the same scan with both barriers per step.
// reverseTilesDivergentBarrier a tile loop whose barriers sit inside a
// guard that diverges on the partial tile.
// Synccheck's target. Day 14's material.
// reverseTilesFixed the guard covers the loads and stores, the
// barriers stand outside it.
//
// The first argument picks which pair runs: `buggy`, `fixed`, or `all` (the
// default). One binary per verdict is what lets a sanitizer transcript name
// one kernel without the other pair's noise in it.
//
// The buggy kernels are reported and never gated. A data race and a divergent
// barrier are both undefined, so a check demanding a wrong answer from them
// would assert that undefined behaviour is dependable. Day 31's evidence has
// the racy scan producing the *right* answer at 32 threads a block, and day
// 27's racy reduction passed 100 runs out of 100. The gates are on the two
// fixed kernels, exact at every block size in the sweep.
//
// Every element is in[i] = i, so every tile prefix is a whole number under
// 2^24 (largest tile sum here is about 8.5e6 at 1024 threads) and a float
// holds it exactly. Comparisons are != with no tolerance: a mismatch means
// the wrong elements were combined, never rounding. All values are distinct,
// so a stale read from the previous tile cannot masquerade as correct.
//
// Nothing is timed. This day is about ordering; the sanitizer serializes and
// replays enough that any clock here would measure the tool.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o racecheck racecheck.cu
// Run: ./racecheck [buggy|fixed|all]
//
#include <cstdio>
#include <cstdlib>
#include <cstring>
#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. main() reads the device's warpSize
// back and fails if it disagrees, because the severity split racecheck is
// predicted to print (WARNING inside a warp, ERROR across warps) is read
// against this number.
constexpr int kWarpSize = 32;
// The shared tiles are sized from this, not from blockDim.x, so one build
// serves every block size the exercise tries, 32 included.
constexpr int kMaxThreadsPerBlock = 1024;
// The block size the lesson's prose quotes; the sweep also runs one warp.
constexpr int kThreadsPerBlock = 256; // 8 warps
// 8 * 1024 + 611. 611 is 13 x 47, so no block size in the sweep divides the
// total and the last tile is always partial, cut mid-warp (99 valid threads
// at 256, 3 at 32). The partial tile is what makes the bad barrier diverge.
// Small on purpose: racecheck instruments every shared access and a day 31
// sized input would turn a two second check into minutes.
constexpr size_t kElems = 8ull * 1024ull + 611ull;
// Tiles each block walks in the reverse kernels. At both sweep sizes the
// last block owns more than one tile, so when it reaches the partial tile
// its shared array holds the previous tile's values, and the buggy kernel's
// stale reads are stale rather than uninitialized.
constexpr size_t kTilesPerBlock = 4;
constexpr int kSweep = 2;
constexpr int kBlockSizes[kSweep] = {32, kThreadsPerBlock};
static_assert(kThreadsPerBlock % kWarpSize == 0,
"block size must be a whole number of warps");
static_assert((kThreadsPerBlock & (kThreadsPerBlock - 1)) == 0,
"the scan doubles its offset, so the block size must be a "
"power of two");
static_assert(kThreadsPerBlock <= kMaxThreadsPerBlock,
"the shared tiles are sized from kMaxThreadsPerBlock");
static_assert(kMaxThreadsPerBlock * sizeof(float) <= 49152,
"one tile for the largest block must fit the 48 KiB a block "
"gets by default on a Tesla T4");
// DELIBERATE BUG, day 62's first target. This is day 31's in-place
// Hillis-Steele tile scan (scanInclusiveInPlace in code/day31-scan-1),
// kept with its bug: one barrier per step where the step needs two.
//
// Thread `tid` reads tile[tid - offset] in the same step in which thread
// `tid - offset` writes tile[tid - offset]; the only barrier sits below
// both, so the read and the write are unordered. The single line inside
// the `if` performs both accesses, which is the line racecheck's report
// is predicted to name. Day 31 measured this kernel producing the right
// answer at 32 threads a block, so the bare-run table below can show a
// zero on this row while the tool still reports the hazard.
//
// Launch assumption: blockDim.x is a power of two, at most
// kMaxThreadsPerBlock, and the grid covers ceil(n / blockDim.x) blocks.
// snippet: racy-step
__global__ void scanTileRacy(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;
__syncthreads();
for (unsigned int offset = 1; offset < blockDim.x; offset *= 2) {
if (tid >= offset) {
tile[tid] += tile[tid - offset]; // read and write, unordered
}
__syncthreads();
}
if (i < n) {
out[i] = tile[tid];
}
}
// end snippet
// The fix: split each step into a read phase and a write phase, with the
// barrier day 31's kernel was missing between them. Every thread reads its
// left neighbour into a register, the whole block passes a barrier, and only
// then does anyone write. The second barrier at the bottom of the loop keeps
// step s+1's reads off step s's writes, which is the job the racy version's
// only barrier was already doing.
__global__ void scanTileFixed(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;
__syncthreads();
// snippet: fixed-step
for (unsigned int offset = 1; offset < blockDim.x; offset *= 2) {
float addend = 0.0f;
if (tid >= offset) {
addend = tile[tid - offset];
}
__syncthreads(); // every read lands before any write below
if (tid >= offset) {
tile[tid] += addend;
}
__syncthreads(); // every write lands before the next step reads
}
// end snippet
if (i < n) {
out[i] = tile[tid];
}
}
// DELIBERATE BUG, day 62's second target. Each block reverses
// kTilesPerBlock consecutive tiles through shared memory, day 14's
// load-barrier-use cycle in a loop, and the guard has swallowed the
// barriers. On every full tile the condition is uniform and nothing is
// wrong. On the last, partial tile the threads past the end skip both
// barriers, so the block arrives at a barrier with divergent threads,
// which is what synccheck exists to report.
//
// Why this does not hang here: the partial tile is the last iteration of
// its block's loop, so the threads that skip the barriers fall out of the
// loop and exit, and __syncthreads() releases when every non-exited
// thread has arrived. That is a property of this loop's geometry, not a
// defence of the pattern; the condition is still not uniform across the
// block, which the programming guide requires.
//
// The answer goes wrong too: the in-range threads read tile cells whose
// owners loaded nothing this iteration, so those reads return the
// previous tile's values. Stale, not racy: nobody writes those cells
// concurrently, so racecheck is predicted to stay silent about this
// kernel while synccheck names it. Two tools, two different bugs.
//
// Launch assumption: gridDim.x covers ceil(numTiles / kTilesPerBlock).
__global__ void reverseTilesDivergentBarrier(const float* __restrict__ in,
float* __restrict__ out, size_t n,
size_t numTiles) {
__shared__ float tile[kMaxThreadsPerBlock];
const unsigned int tid = threadIdx.x;
const unsigned int width = blockDim.x;
const size_t first = static_cast<size_t>(blockIdx.x) * kTilesPerBlock;
const size_t last =
(first + kTilesPerBlock < numTiles) ? first + kTilesPerBlock : numTiles;
// snippet: divergent-barrier
for (size_t t = first; t < last; ++t) {
const size_t i = t * width + tid;
if (i < n) {
tile[tid] = in[i];
__syncthreads(); // divergent on the partial tile
out[i] = tile[width - 1u - tid];
__syncthreads(); // divergent again
}
}
// end snippet
}
// The fix is day 14's rule applied inside a loop: guard the loads and the
// stores, never the barrier. Every thread of the block reaches both barriers
// on every iteration, out-of-range threads load the 0.0f identity, and the
// second barrier keeps iteration t+1's loads off iteration t's reads.
__global__ void reverseTilesFixed(const float* __restrict__ in,
float* __restrict__ out, size_t n,
size_t numTiles) {
__shared__ float tile[kMaxThreadsPerBlock];
const unsigned int tid = threadIdx.x;
const unsigned int width = blockDim.x;
const size_t first = static_cast<size_t>(blockIdx.x) * kTilesPerBlock;
const size_t last =
(first + kTilesPerBlock < numTiles) ? first + kTilesPerBlock : numTiles;
// snippet: fixed-barrier
for (size_t t = first; t < last; ++t) {
const size_t i = t * width + tid;
tile[tid] = (i < n) ? in[i] : 0.0f;
__syncthreads(); // every thread, every iteration
if (i < n) {
out[i] = tile[width - 1u - tid];
}
__syncthreads(); // the tile is rewritten next iteration
}
// end snippet
}
// CPU reference for the tile scan: inclusive, restarting at every tile
// boundary. Accumulates in double where the kernel accumulates in float; on
// this input both are exact, so the cast back loses nothing.
static void scanTilesCpu(const float* in, float* out, size_t n, size_t tile) {
for (size_t base = 0; base < n; base += tile) {
const size_t end = (base + tile < n) ? base + tile : n;
double running = 0.0;
for (size_t i = base; i < end; ++i) {
running += static_cast<double>(in[i]);
out[i] = static_cast<float>(running);
}
}
}
// CPU reference for the tile reverse. Element j of a tile takes element
// width-1-j of the same tile, and a source past the end of the array is the
// 0.0f the fixed kernel's padded load supplies.
static void reverseTilesCpu(const float* in, float* out, size_t n,
size_t width) {
for (size_t base = 0; base < n; base += width) {
for (size_t j = 0; j < width && base + j < n; ++j) {
const size_t src = base + (width - 1 - j);
out[base + j] = (src < n) ? in[src] : 0.0f;
}
}
}
// Counts the elements where got and want differ and reports the first one.
// Exact, no tolerance: see the header note on why every value here is a
// whole number a float represents exactly.
static size_t countMismatches(const float* got, const float* want, size_t n,
size_t* firstBad) {
size_t bad = 0;
*firstBad = n;
for (size_t i = 0; i < n; ++i) {
if (got[i] != want[i]) {
if (bad == 0) {
*firstBad = i;
}
++bad;
}
}
return bad;
}
int main(int argc, char** argv) {
// Which pair of kernels runs. One verdict per sanitizer invocation is
// what keeps each transcript about one bug.
bool runBuggy = true;
bool runFixed = true;
if (argc > 1) {
if (std::strcmp(argv[1], "buggy") == 0) {
runFixed = false;
} else if (std::strcmp(argv[1], "fixed") == 0) {
runBuggy = false;
} else if (std::strcmp(argv[1], "all") != 0) {
std::fprintf(stderr, "usage: %s [buggy|fixed|all]\n", argv[0]);
return EXIT_FAILURE;
}
}
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
std::printf("GPU: %s (compute capability %d.%d), warp size %d\n", prop.name,
prop.major, prop.minor, prop.warpSize);
std::printf("mode: %s%s%s\n", runBuggy ? "buggy" : "",
(runBuggy && runFixed) ? "+" : "", runFixed ? "fixed" : "");
// Every failure below records itself and falls through to the one
// cleanup block at the bottom, so no path returns with device memory
// allocated.
int failures = 0;
if (prop.warpSize != kWarpSize) {
std::fprintf(stderr,
"this card reports warp size %d; the predicted severity "
"split is read against %d\n",
prop.warpSize, kWarpSize);
++failures;
}
const size_t bytes = kElems * sizeof(float);
std::vector<float> h_in(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = static_cast<float>(i);
}
std::vector<float> h_out(kElems);
std::vector<float> h_want(kElems);
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));
std::printf(
"\n%zu elements, in[i] = i, one tile per scan block, %zu "
"tiles per reverse block\n",
kElems, kTilesPerBlock);
std::printf("%7s %-32s %11s %11s %6s\n", "threads", "kernel", "mismatches",
"first bad", "gated");
std::printf("%7s %-32s %11s %11s %6s\n", "-------",
"--------------------------------", "----------", "---------",
"-----");
for (int s = 0; s < kSweep; ++s) {
const int threads = kBlockSizes[s];
if (threads > prop.maxThreadsPerBlock) {
std::fprintf(stderr,
"this card caps a block at %d threads; the %d row "
"cannot run\n",
prop.maxThreadsPerBlock, threads);
++failures;
continue;
}
const size_t width = static_cast<size_t>(threads);
const size_t numTiles = (kElems + width - 1) / width;
const int scanBlocks = static_cast<int>(numTiles);
const int reverseBlocks =
static_cast<int>((numTiles + kTilesPerBlock - 1) / kTilesPerBlock);
// Four launches, table order: racy scan, fixed scan, divergent
// reverse, fixed reverse. `gated` marks the two whose mismatch
// count is a build verdict rather than a report.
for (int kernel = 0; kernel < 4; ++kernel) {
const bool buggy = (kernel == 0 || kernel == 2);
if ((buggy && !runBuggy) || (!buggy && !runFixed)) {
continue;
}
const bool isScan = (kernel < 2);
const char* name = nullptr;
if (isScan) {
scanTilesCpu(h_in.data(), h_want.data(), kElems, width);
} else {
reverseTilesCpu(h_in.data(), h_want.data(), kElems, width);
}
switch (kernel) {
case 0:
name = "scanTileRacy";
scanTileRacy<<<scanBlocks, threads>>>(d_in, d_out, kElems);
break;
case 1:
name = "scanTileFixed";
scanTileFixed<<<scanBlocks, threads>>>(d_in, d_out, kElems);
break;
case 2:
name = "reverseTilesDivergentBarrier";
reverseTilesDivergentBarrier<<<reverseBlocks, threads>>>(
d_in, d_out, kElems, numTiles);
break;
default:
name = "reverseTilesFixed";
reverseTilesFixed<<<reverseBlocks, threads>>>(
d_in, d_out, kElems, numTiles);
break;
}
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
size_t firstBad = kElems;
const size_t bad =
countMismatches(h_out.data(), h_want.data(), kElems, &firstBad);
char firstText[24];
if (bad == 0) {
std::snprintf(firstText, sizeof(firstText), "%s", "-");
} else {
std::snprintf(firstText, sizeof(firstText), "%zu", firstBad);
}
std::printf("%7d %-32s %11zu %11s %6s\n", threads, name, bad,
firstText, buggy ? "no" : "yes");
// Only the fixed kernels are gated. Demanding a wrong answer
// from a race or a divergent barrier would gate on undefined
// behaviour, and both buggy rows are allowed to read zero.
if (!buggy && bad != 0) {
std::fprintf(stderr,
"%s at %d threads: %zu of %zu elements wrong, "
"first at %zu, got %.9g want %.9g\n",
name, threads, bad, kElems, firstBad,
static_cast<double>(h_out[firstBad]),
static_cast<double>(h_want[firstBad]));
++failures;
}
}
}
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
std::printf(
"\na zero in a buggy row is a report, not a verdict; the "
"verdict is the sanitizer's\n");
if (failures != 0) {
std::fprintf(stderr, "%d check(s) failed\n", failures);
return EXIT_FAILURE;
}
if (runFixed) {
std::printf(
"both fixed kernels matched the reference at every "
"block size\n");
}
return EXIT_SUCCESS;
}