← Chapter 23: What syncthreads protects Chapter 25: What makes CUDA code "correct" →
date 12 / jul / 26
mtwtfss
SYNCHRONIZATION & HAZARDS — Chapter 24 index home

Why race conditions happen

*the lottery*

Launch two threads. Both run the same one line, increment a counter. The counter starts at zero. What is the final value? It feels obvious that it is two, one for each thread. But it might be one. Two increments went in and only one came out, and nothing crashed, nothing warned you. Your program is just quietly wrong. This is a race condition, where correctness depends on the unpredictable timing of who runs when. Understanding it is the difference between a kernel you can trust and a kernel that lies to you on some runs and not others.

THE SIMPLE PROGRAM THAT LIES TO YOUCounter = 0Launch two threads.Both run the SAME line:counter = counter + 1;What should the final value be?2What you might get:1...and nothing crashes.Nothing warns you.Just quietly wrong.A race condition: correctness hinges on the unpredictable timingof who runs when - the result can change every single run.

The Trap: Read-Modify-Write

Start with why incrementing a counter is dangerous, because it looks harmless. The trap is that increment is not one operation. It is three. First the thread reads the counter's current value out of memory into a private register. Then it adds one to that register. Then it writes the register back to memory. Read, modify, write. Three separate steps with real gaps between them, and those gaps are where the disaster lives.

WHY INCREMENT IS DANGEROUS: 3 OPS, NOT 11. READRead the counterfrom memoryinto a register.Memory0Register02. MODIFYAdd onein the register.0+113. WRITEWrite the registerback to memory.Register1Memory1Read, modify, write - three steps, with real gaps between them.

Walk two threads through it in the unlucky order. Thread A reads the counter and sees zero. Before thread A writes anything back, thread B also reads the counter, and it also sees zero, because A hasn't written yet. Now both threads hold zero in their private registers. Both add one, both compute one. Thread A writes one back. Then thread B writes one back, right on top of A's result.

THE UNLUCKY INTERLEAVING (FINAL = 1)StepThread AThread BMemory1Read → 002Read → 003Add 1 → 1Add 1 → 104Write 115Write 11RESULT: one increment is lost.Both wrote 1 over the same old value.

↳ The counter ends at one. Thread B never saw A's update, because it read the value before A's write had landed. One of the two increments simply evaporated. Scale this to ten thousand threads and you don't get ten thousand, you get some smaller, random number that changes every run.

Is it a bug, or the hardware?

Now the natural question. Is this just a bug in my code, or is the hardware actually working against me? It is the hardware, by design, for three concrete reasons.

  1. Scheduling is a lottery: The order threads run in is not predictable. The scheduler on each processing unit runs whichever group of threads happens to have its data ready, and swaps groups in and out as they stall on memory. You cannot know which thread reaches the counter first. The lottery is not an accident, it is the normal way the machine operates.
  2. Multiple steps: The read-modify-write is genuinely multiple steps, so there is always a window between reading and writing where another thread can slip in and read the same stale value.
  3. Delayed visibility: A write does not appear to everyone instantly. It takes time to travel through the layers of memory and caches before other threads can see it. So a reader can grab the old value from its own cache before the new one has propagated across. Even threads close together do not have a perfectly shared, instant view of memory.
WHY THE HARDWARE WORKS AGAINST YOU (BY DESIGN)1. UNPREDICTABLEORDERScheduler runs whoeveris ready - you can't knowwho goes first.SchedulerWarpAWarpBWarpCAny order!2. RMW ISMULTI-STEPAlways a windowbetween read and writefor another threadto slip in.ReadModifyWrite----THE GAP3. WRITES ARENOT INSTANTWrites travel throughcaches & memory beforeothers can see them.Core A (write)CacheMemory...Core B (read)timeThe lottery is not a bug - it's the normal way the machine runs.

Why the barrier fails here

So how do you fix it? And here is the most important and most misunderstood point in all of this. The barrier you use to make threads wait for each other does not fix this race.

It is tempting to think, just make everyone pause between reading and writing. But watch it fail. Every thread reads the counter before the pause, so they all see the same value. Then they all wait at the pause. Then they all write their incremented value, and the last one to write still stomps on everyone else.

THE BARRIER DOES NOT FIX THIS RACEIt only orders the phases, not the single operation.BAD: a barrier between read & write? NOPE.All threads read first (all see 0):Read 0Read 0Read 0...__syncthreads();...then all write (all write 1):Write 1Write 1Write 1...Last writer wins. Others are lost.The barrier ordered the phases,not the single update.Recall (Ch23): a barrier givesSTOP + FLUSH - order + visibility.Great for hand-off:one writer, many readers.NOT for many writers.A barrier orders WHEN things happen - it does not makethe three-step increment indivisible.

↳ The pause synchronized the phases, everyone read then everyone wrote, but it did nothing to protect the single update from being overwritten. A barrier orders when things happen, it does not make the three-step increment indivisible.

The Real Fix: Atomics

What actually fixes it is an atomic operation. An atomic increment performs the read, the modify, and the write as one single indivisible hardware transaction.

While one thread is doing it, the hardware locks that memory location so no other thread can read or write it until the whole read-modify-write is complete. The threads are forced to take turns on that one spot. So ten thousand atomic increments correctly produce ten thousand, because none of them can slip into the gap of another, there is no gap.

THE REAL FIX: ATOMIC OPERATIONSWithout Atomic (broken)Read 0Read 0Read 0each: +1, then Write 11Lost updates!With Atomic (correct)ATOMIC ADD (counter, 1)Hardware locks the location.Read + Modify + Write asONE indivisible step.0 → 1 → 2 → 3 → ... → NN = number of threadsAtomics serialize access to one spot: slower, but correct.Slow and correct beats fast and wrong.

↳ The cost is that it is slower than a plain increment, because it serializes access. Threads line up for that location instead of charging in parallel. But slow and correct beats fast and wrong.

It is worth being precise about when a barrier is the right tool and when it is not, because both exist for real reasons. A barrier is perfect for a single-writer, many-reader pattern. One thread computes a value and writes it, everyone hits the barrier so that write is guaranteed done and visible, then everyone safely reads it. No race, because only one thread ever writes. The barrier protects readers from seeing a half-finished write.

Atomics are for the opposite situation, many threads all updating the same location, where you need each update to survive. Use the barrier to separate a writing phase from a reading phase. Use atomics when many writers hit one spot. They solve different problems and are not interchangeable.

WHEN TO USE WHATUse a BARRIERsingle-writer, many-reader1 writer computes...writes →SharedValue__syncthreads();...many readers consume:• only one thread writes• write is done & visible before reads• readers never see half-writesUse ATOMICSmany-writer, single locationatomicAdd(counter, 1)each update survives.• many threads, one location• each update is indivisible• no updates are lost• higher cost (serialization)Different problems, different tools - they are not interchangeable.

The Warp-Level Sync

There is also a finer-grained cousin of the block barrier worth knowing. The full barrier makes every thread in the whole block wait, and it lives in the block's shared memory hardware. But sometimes you only need the 32 threads of a single warp to agree, and for that there is a lighter warp-level sync.

And this raises a fair question. Don't the 32 threads of a warp already run in perfect lockstep, so why sync them at all? They used to, but on modern hardware that is no longer guaranteed.

Newer architectures let threads within the same warp drift to different points in the program and reconverge, so you can no longer assume that if one thread in a warp wrote something, its neighbor has already seen it. The warp-level sync makes the 32 threads line up and agree on memory, and it is cheaper than the full block barrier because it only involves one warp's worth of hardware, not the whole block. If all you need is agreement inside a warp, reach for the lighter one.

WARP-LEVEL vs BLOCK-LEVEL SYNCBLOCK-LEVEL(__syncthreads)All threads in the block.Whole Block• involves all threads• higher cost (more HW)• needed when threadsoutside the warp must agreeWARP-LEVEL(__syncwarp(mask))Only the warp's 32 threads.One Warp (32)Why needed?Old GPUs: warp = strict lockstep.New GPUs: threads can diverge& reconverge - visibility not assured.__syncwarp() forces agreement,cheaper than a full block barrier.Good RuleUse the smallestsync that solvesyour problem.warp sync <block syncReach for the smallest synchronization that solves your problem.

↳ The General Takeaway

So the whole picture. A race condition is when your answer depends on the accidental timing of threads, and it strikes hardest on read-modify-write, because that trio of steps has gaps that another thread can invade, leaving updates to vanish silently.

The hardware makes this the default through unpredictable scheduling, multi-step updates, and delayed propagation. A barrier does not save you here, it only orders phases, it cannot make an increment indivisible. Atomics are the real fix for many-writer updates, locking the location so each update survives at the cost of serializing.

Shared memory on a GPU is a commons with no automatic protection.
The hardware will not stop you from writing wrong code, it will just quietly give you a different wrong answer every time.