Race Conditions Don't Disappear at Low Precision — They Hide

The challenge

Quantized training and inference push real numerical work into fp16, bf16, fp8, and fp4. That migration is usually framed as a question of accuracy — how much precision can a model give up? A quieter and more dangerous question goes unasked: does lowering precision change how concurrency bugs behave?

A data race is undefined behaviour regardless of the datatype [8]. But whether an engineer ever sees it — in a unit test, in a validation metric, in a customer incident — depends on whether the corrupted bytes happen to differ from the correct ones. At low precision, neighbouring values collapse onto the same bit pattern [5], and a race that would be obvious in fp32 can return the "right" answer by accident. Layer on a hardware generation change, and the same source code can flip from silently-correct to materially-wrong.

This investigation set out to answer three concrete questions a team shipping quantized GPU code actually needs answered:

1. Do races behave differently across precision and GPU architecture?
2. Do atomics always protect you — and what does synchronization cost?
3. Can compute-sanitizer catch these defects on the newest Blackwell Ultra silicon?

The approach

Rather than hunt for races in production code, the study manufactures them: seven minimal race patterns (A–G), each paired with a correct "safe twin" so the only variable is the missing synchronization. Each pattern was swept across six precisions, three NVIDIA GPU generations spanning Turing to Blackwell Ultra, and mirrored on the CPU in OpenMP to contrast a barrier-rich language model. Two instrumentation layers sat on top: cycle-accurate cost microbenchmarks (clock64()), and the full compute-sanitizer tool suite (racecheck, synccheck, memcheck, initcheck) [3]. The synchronization semantics under test — __syncthreads, atomics, warp-shuffle masks, and memory fences — follow the CUDA memory consistency model [1], whose weakly-ordered behaviour on NVIDIA GPUs is documented experimentally by Alglave et al. [2].

The patterns deliberately span the hazard taxonomy: global read-modify-write without atomics, shared-memory ordering, warp-shuffle masks, sub-word (fp8/fp4) byte and nibble collisions, reduction synchronization, cross-stream memory fences, and tensor-core accumulation.

What we found

Low precision hides the bug

Pattern B — a shared-memory read across warps with a missing __syncthreads() — exposed the core risk directly. On a B300 at one million elements:

The race is identical in all four cases. What changes is that low-precision neighbours round to the same bytes, so a stale read usually returns a value indistinguishable from the correct one. The defect ships silently inside the exact datatypes quantized ML depends on — including the fp8 e4m3/e5m2 encodings now standard for deep learning [6].

The same code, 119× more dangerous on newer hardware

Holding the kernel and workload fixed and varying only the GPU:

A 119× swing with zero source changes. Blackwell's more aggressive warp scheduler keeps more work in flight and widens the window between an unsynchronized write and a neighbouring read — precisely the weakly-ordered behaviour that GPU memory-model studies warn against relying on [2]. "It passed on our A6000" is not a correctness argument.

Tensor cores: a single non-atomic store corrupts everything

Pattern G models split-K matrix multiply, where several blocks each compute a partial product and accumulate into the same output tile. The natural-looking wmma::store_matrix_sync is not atomic — and there is no atomic-add for tensor-core matrix fragments at all. The result was 100% of tiles corrupted: blocks overwrite instead of sum. The only correct design is per-block workspace plus an explicit reduction — which is precisely why production libraries (cuBLAS, CUTLASS) structure split-K that way [4]. The study demonstrates it is mandatory, not an optimization.

Atomics are necessary, not sufficient

Four documented cases where the atomic is individually correct yet the program still races: ABA, composability (a load + atomicSub pair drove a counter to −64,448) [7], memory ordering [1, 8], and sub-word types with no atomic support. Atomics protect a single-variable read-modify-write — and nothing larger.

Tooling lags the hardware — verify, don't assume

compute-sanitizer could not instrument the B300 under CUDA 12.8 or 12.9.2 (both report "Device not supported") [3, 9]; it required CUDA 13.3. Once working, racecheck reported the identical 2,097,160 hazards on all three GPUs — because it counts hazards from the static access pattern, not runtime timing. That is the synthesis of the whole study in one number: the hazard exists identically everywhere; how often it bites is architecture- and precision-dependent. A race that "never reproduces" on a Turing dev box is the same defect quietly corrupting a quarter of reads in Blackwell production.

The cost of doing it right

Fixing these defects is cheap if you pick the right primitive. Measured on the B300:

The engineering guidance that falls out is consistent across GPU and CPU: privatize and reduce rather than hammering a contended shared atomic, and reach for mutual exclusion only when a genuine multi-step invariant demands it.

Why it matters for clients

This is the kind of defect that passes every test and ships anyway — then surfaces as an irreproducible accuracy regression after a hardware refresh or a precision change. The deliverables here are directly reusable: a portable race stress-harness, a cross-architecture correctness baseline, a sync-primitive cost model, and a reproducible compute-sanitizer workflow (including the toolkit version gotchas) that a team can run against its own kernels.

The full source, per-architecture data, and a presentation deck are linked below.

References

1. NVIDIA. CUDA C++ Programming Guide — memory consistency model, __syncthreads(), atomic functions, warp-matrix (WMMA) functions, __threadfence(). https://docs.nvidia.com/cuda/cuda-c-programming-guide/
2. J. Alglave, M. Batty, A. F. Donaldson, G. Gopalakrishnan, J. Ketema, D. Poetzl, T. Sorensen, J. Wickerson. "GPU Concurrency: Weak Behaviours and Programming Assumptions." ASPLOS, 2015. https://dl.acm.org/doi/10.1145/2694344.2694391
3. NVIDIA. Compute Sanitizer User Manual — racecheck, synccheck, memcheck, initcheck; supported devices. https://docs.nvidia.com/compute-sanitizer/
4. NVIDIA. CUTLASS — efficient GEMM and split-K parallel reduction (per-partition workspace + reduction). https://github.com/NVIDIA/cutlass
5. D. Goldberg. "What Every Computer Scientist Should Know About Floating-Point Arithmetic." ACM Computing Surveys, 23(1), 1991. https://dl.acm.org/doi/10.1145/103162.103163
6. P. Micikevicius, et al. "FP8 Formats for Deep Learning." arXiv:2209.05433, 2022. https://arxiv.org/abs/2209.05433
7. M. Herlihy, N. Shavit. The Art of Multiprocessor Programming, 2nd ed. Morgan Kaufmann, 2020 — ABA problem, lock-free composability.
8. S. V. Adve, H.-J. Boehm. "Memory Models: A Case for Rethinking Parallel Languages and Hardware." Communications of the ACM, 53(8), 2010 — data races as undefined behaviour, memory ordering. https://dl.acm.org/doi/10.1145/1787234.1787255
9. NVIDIA. Blackwell Compatibility Guide — compute capability 10.3 (sm_103), B300 toolchain support. https://docs.nvidia.com/cuda/blackwell-compatibility-guide/