100 DAYS OF GPU Course preview
Day 62Module 7
in-technical-review

Finding races with racecheck and synccheck

Day 27 ran a racy reduction one hundred times and it produced the right answer one hundred times. Day 31 ran a racy scan at 32 threads a block and its mismatch count was zero. Both kernels are wrong, but their output tests kept passing.

Day 14 introduced a tool that checks ordering without relying on the output. This page uses compute-sanitizer to trace a racecheck hazard to its source line. It also explains why synccheck stays silent on a guarded barrier and states what neither tool can see.

What each tool checks, and what it cannot see

compute-sanitizer is one binary with four tools behind a --tool flag. Day 61 used memcheck, the default, which catches bad addresses. Today's two catch bad ordering, and they split the work cleanly.

racecheck watches shared memory. The manual defines it as follows: "The racecheck tool is a run time shared memory data access hazard detector" (https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html , checked 2026-09-01). It instruments every shared access and records, for each cell, whether two accesses from different threads arrive with no barrier between them and at least one is a write.

The tool names three hazard classes by observed order: RAW (a read after an unordered write), WAR (a write after an unordered read), and WAW (two unordered writes). Each hazard has a severity. ERROR marks an error, while WARNING marks "hazards due to warp level programming that make the assumption that threads are proceeding in groups" (same manual, same date): two threads of one warp touching one cell, which old code assumed lockstep made safe and day 14 showed the scheduler stopped promising on this hardware generation.

synccheck watches the barrier itself. The manual again: it "can identify whether a CUDA application is correctly using synchronization primitives, specifically __syncthreads() and __syncwarp() intrinsics and their Cooperative Groups API counterparts". Where racecheck asks whether the data needed a barrier it did not get, synccheck asks whether the barrier you did write is legal: every __syncthreads() must be reached by the whole block, so a barrier under a condition that splits the block is a natural candidate for a barrier error with divergent threads, the manual's categories being "Divergent thread(s) in block" and "Divergent thread(s) in warp". The measured guarded-barrier kernel on this page is the important counterexample: synccheck reports neither category, under both compute-sanitizer 12.6 and 13.0.85.

The same definition states racecheck's limit: "Currently, this tool only supports detecting accesses to on-chip shared memory." A race through global memory, which is exactly what day 27's reduction had, produces no hazard, no warning, nothing.

The toolkit has no global-memory race checker. For that class of bug, use the memory model: name the store, the read, and the operation that orders them.

Diagram: two bugs, two tools. Three bands, each a slice of one block's timeline with shared memory as a strip below the threads. Band 1, the racy scan step: thread 2 writes tile cell 2 while thread 3 reads it, no barrier between the two arrows. Caption "one line, two unordered accesses: a hazard, whichever value the read got." Band 2, the fixed step: the same two arrows with a barrier bar between the read phase and the write phase. Caption "two barriers per step, zero hazards." Band 3, the partial tile: 32 lanes reach a guard, 3 pass it to the barrier bar, 29 turn away before it. Caption "3 of 32 lanes arrive: a divergent barrier." Alt text: "A race is two unordered arrows to one shared cell and a barrier bar between them fixes it. A divergent barrier is three lanes of thirty-two arriving while the rest leave."

A right answer is not a clean kernel

Correctness testing does not catch a race as reliably as an off-by-one error. A single-threaded bug often produces the same wrong bytes on each run, while a race is a missing order. The scheduler may choose the expected order on a given launch.

Day 27's kernel did so for one hundred launches in a row. Day 31's in-place scan did so whenever the block contained one warp.

racecheck ignores the output and checks whether accesses were ordered. It therefore reports the same hazards when a run produces the expected answer. A zero mismatch count describes that run but does not prove the kernel has no race.

One file, two bugs, a switch

Full program in code/day62-racecheck/racecheck.cu. It plants both bugs on purpose, labels them as planted in the comments, and ships each fix in the same file. The habits on top are what make the transcripts predictable.

One mode per transcript. The first argument picks buggy, fixed or all, so a sanitizer run over the bugs is not polluted by the fixes and a run over the fixes can be predicted clean.

The buggy rows are never gated. Both bugs are undefined behaviour, so the gates (real branches returning EXIT_FAILURE) sit on the two fixed kernels only, exact against a CPU reference at both block sizes.

Nothing is timed. racecheck instruments every shared access; any clock in this program would measure the tool.

The first bug keeps day 31's in-place tile scan with the barrier it was missing. One line reads a neighbour's cell and writes its own, and the only barrier sits below the 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];
    }
}

The fix splits every step into a read phase and a write phase with the missing barrier between them:

    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
    }

The second bug is day 14's material inside a tile loop: each block reverses four tiles through shared memory, and the guard has swallowed the barriers. On every full tile the condition is uniform and nothing diverges. On the last, partial tile, the threads past the end of the array skip both barriers:

    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
        }
    }

That kernel does not hang here because the partial tile is the final loop iteration for its block. The threads that skip the guard exit before the remaining threads reach the barrier. This input avoids a hang but does not make the code safe.

The guide requires the condition to be uniform across the block, the answer comes out wrong (the in-range threads read cells nobody loaded this lap), and moving the partial tile anywhere but the last lap of the loop turns the same pattern into a hang. The fix is day 14's rule with a second barrier for the loop: guard the loads and stores, never the barriers, and let the out-of-range threads carry the 0.0f identity.

    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
    }

Results

Run on the project's Tesla T4 (driver 595.84, CUDA 12.6 V12.6.85) on 2026-09-01, compute-sanitizer 12.6. Six transcripts in code/day62-racecheck/evidence/. Three predictions held and one died, and the one that died is the most useful line on the page.

CUDA 13.0 with compute-sanitizer 13.0.85 preserved those verdicts. Analysis mode again displayed 1 error and 1 warning for scanTileRacy, now at source line 128; hazard mode again hit the 100-report display limit; both fixed runs were clean; and synccheck remained silent on the buggy divergent barrier. The aggregated race hazard counts changed, as instrumentation and scheduling allow, but the bare-run mismatch table was identical.

Held: racecheck names scanTileRacy and nothing else, at one line. Analysis mode reported two hazards, RACECHECK SUMMARY: 2 hazards displayed (1 error, 1 warning), both between a read and a write inside scanTileRacy at racecheck.cu:131. The severity split held exactly as predicted: the WARNING carries 108,192 hazards and the ERROR carries 213,492. Hazard mode expands the same picture to 100 hazards displayed (74168 errors, 248332 warnings), the 100 being the tool's print limit rather than a count of anything.

Held: both tools report zero on fixed. RACECHECK SUMMARY: 0 hazards displayed (0 errors, 0 warnings), and synccheck's ERROR SUMMARY: 0 errors. The fixes are fixes.

Died: synccheck reported ERROR SUMMARY: 0 errors on the buggy build. It found nothing in reverseTilesDivergentBarrier, on either block size, while the kernel's own output is visibly wrong on the same run: 3 mismatches at 32 threads and 99 at 256. The prediction said at least one barrier error per block size. It got none.

This result narrows what synccheck proves. The tool detects illegal barrier usage, threads of a warp arriving at different barriers, and this kernel does something legal-looking and still wrong: every thread that reaches a barrier reaches the same one, but the threads the guard excluded never arrive.

The block's remaining threads pass a barrier that was supposed to wait for them. The tool has no complaint because no warp diverged at a barrier; the data is wrong because the barrier did not mean what the code assumed. Day 14 taught that an early return before a barrier does not hang.

This is the same case with an added failure: it does not hang or trigger the tool, but it produces wrong answers.

Held: each tool missed the other defect. racecheck said nothing about the barrier kernel, and synccheck said nothing about the scan race. A clean report covers only the defect class that tool checks.

Kernel, buggy mode Bare-run mismatches (32 / 256 threads) racecheck synccheck
scanTileRacy 0 / 0 1 error, 1 warning at line 131 silent
reverseTilesDivergentBarrier 3 / 99 silent silent

In the plain run, scanTileRacy produced zero mismatches at both block sizes, so the output did not expose the race. Under racecheck it produced 5,781 mismatches at 256 threads, and under synccheck 7,138. Instrumentation changed the timing enough to expose wrong answers.

A passing output check does not prove that a racy kernel is correct.

Run it yourself

Run this on a CUDA-capable GPU with Compute Sanitizer installed. The tool reads no performance counters, so day 42's permission error does not apply. The README contains the build line and five commands.

NVIDIA's dev-tools notebook includes a Compute Sanitizer example: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/cuda-cpp/notebooks/03.02-Kernels/03.02.04-Dev-Tools.ipynb (checked 2026-09-01).

Exercise

Fix both bugs in your copy of racecheck.cu, without looking at the fixed kernels, until both tools report zero in buggy mode too. Then point racecheck at day 27's memory_model binary and explain the silence.

Time: 30 to 40 minutes. Submit: your two edited kernels, the two clean summary lines, and one sentence on what racecheck said about day 27's racy reduction and why.

Check: the harness gates the fixed kernels against a CPU reference, exactly, at 32 and 256 threads a block; a failure names the kernel, the mismatch count and the first bad index, and returns EXIT_FAILURE through a real branch. Your edited kernels must pass that gate and both sanitizer tools. Harness contract at /reference/harness.

Hint 1

For the scan, the report names one line twice. What are the two accesses on that line, which threads make each, and what stands between them? For the loop, ask who is allowed to skip a __syncthreads(), and what the guard is protecting that the barrier is not.

Hint 2

The scan needs its reads finished before any writes start, which no single barrier below the step can order. The loop needs every thread at every barrier, so the load wants an identity value where the array has ended. For day 27: reread racecheck's first sentence in the manual, then say which memory the reduction's partials live in.

Solution

The scan fix is a barrier between reading tile[tid - offset] and writing tile[tid], which in practice means staging the read into a register, or double buffering as day 31 did. The loop fix hoists both barriers out of the guard and pads the load with 0.0f. Racecheck says nothing about day 27, because those partial sums travel through global memory and the tool only watches shared; the silence is a boundary of the tool, not a verdict on the kernel.

A sanitizer's clean summary means no bug of the class it checks, in the code it saw run. Everything outside that class, global races above all, is still yours.

Pitfalls

racecheck reports nothing, so you conclude the kernel is race free. It only watches shared memory, by its own documentation. A global-memory race like day 27's is invisible to it, and no compute-sanitizer tool covers that class. Clean means clean where it looked.

You read WARNING as ignorable. The warp-level severity exists because old warp-synchronous code raced on purpose and mostly won. Day 14 quoted the guide ending that era: sub-warp divergence makes lockstep assumptions invalid on this hardware. Treat a WARNING in code you did not deliberately warp-program as an ERROR with better manners.

Your report shows addresses where this page shows lines. Compile with -lineinfo. It does not change the optimization level, so it belongs in the build line permanently, not just on debugging days.

The mismatch table changes under the sanitizer. Instrumentation changes timing, and a race's outcome is a function of timing. A racy kernel that was wrong bare can come out right under the tool, and the other way round. The hazard report is the stable signal; the mismatch count never was.

You ran compute-sanitizer ./app and expected hazards. The default tool is memcheck. Racing and divergent barriers need --tool racecheck and --tool synccheck, one tool per run.

Your divergent barrier deadlocks on another input. This page's version exits because the divergence lands on the loop's last lap. Shift the partial tile earlier, or add work after the guard, and the threads that skipped the barrier are still alive inside the loop while the rest wait forever. Warp divergence around a barrier is not a performance bug, it is a correctness bug that sometimes hangs.

Go deeper

Next

Day 63 uses cuda-gdb on a kernel that computes a wrong index, stepping one thread while its warp waits, which is the tool for the bugs that are deterministic and still baffling. Day 64 covers printf and device-side assert, the debugging most people try first, and why the output may show events in a different order.