100 DAYS OF GPU Course preview
Day 27Module 3
in-technical-review

The CUDA memory model: fences, scopes and atomics

Someone who had done the reading, on Stack Overflow, 30 votes and 21,796 views:

"I have gone through many forum posts and the NVIDIA documentation, but I couldn't understand what __threadfence() does and how to use it. Could someone explain what the purpose of that intrinsic is?"

https://stackoverflow.com/questions/5232689/what-is-the-purpose-of-the-threadfence-intrinsic-in-cuda (checked 2026-08-30)

The documentation tells you what a fence orders. It cannot tell you which two lines of your kernel need that order. You need a concrete kernel to apply the definition.

This page uses the multi-block reduction from day 24 and day 25, finished in one launch by letting the block that arrives last add the other partial sums. One fence separates the defined version from a version that may read a partial sum before it becomes visible.

What the last block has to be told

Days 24 and 25 stop one step short. A grid of 1,024 blocks turns a large array into 1,024 partial sums in global memory, and something still has to add those. The usual answer is a second launch, because a launch boundary is a device-wide ordering point you get for free.

The other answer is to let the last block do it. Every block stores its partial sum, then atomically adds 1 to a counter and reads the counter's old value. Exactly one block reads gridDim.x - 1, knows the other 1,023 have finished, sums the partials, and writes the answer in the same launch.

NVIDIA ships this method as threadFenceReduction. The whole grid need not be resident because a block may store its partial and retire. Day 28 shows the different residency rule for cooperative groups.

The counter does not order the partial sum store. atomicAdd promises only atomicity. From the guide's atomics section: the legacy atomic functions, meaning every atomic<Op>() this course has used, "have a memory ordering of cuda::std::memory_order_relaxed and are only atomic at a particular thread scope", and "Unlike built-in atomic functions, legacy atomic functions only ensure atomicity and do not introduce synchronization points (fences)." (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html , section 5.4.5, checked 2026-08-30.)

The counter value does not carry or wait for the partial sum. It also does not stop another block from seeing the counter update before the partial sum.

The guide says this about the kernel: "Without a fence between storing the partial sum and incrementing the counter, the counter may increment before the partial sum is stored. This could cause the counter to reach gridDim.x - 1 and allow the last block to start reading partial sums before they are updated in memory."

Blocks can run on different SMs, each with its own L1. The code must ask for the needed order.

Across 1,024 blocks, an atomic ticket may become visible before its preceding partial sum unless __threadfence orders the store before the ticket.

The intuition that an atomic is a barrier

Many programmers expect an atomic operation to synchronize other writes. In C++, std::atomic defaults to sequential consistency, so a store through it publishes earlier writes. CUDA's atomicAdd uses relaxed ordering instead.

It takes no memory order argument. The guide states the order it uses.

The other half of the intuition is that __syncthreads() already covers this. Day 14 is where it stops: both of the barrier's promises end at the block, because the shared memory they exist to order ends there too.

Three fences, one per scope. Each orders your own thread's writes against your own thread's later writes, and the scope names who is guaranteed to see that order. The libcu++ spelling is cuda::atomic_thread_fence(cuda::memory_order_seq_cst, <scope>).

Intrinsic Whose view of your writes it orders Scope
__threadfence_block() "all threads in the calling thread's block" cuda::thread_scope_block
__threadfence() "any thread in the device" cuda::thread_scope_device
__threadfence_system() "all threads in the device, host threads, and all threads in peer devices" cuda::thread_scope_system

The same three scopes name the atomics: atomicAdd is device scope, atomicAdd_block is block scope, atomicAdd_system is system scope, and the suffix is the only thing that says so. Use one outside its scope and you have not written a slower correct program, you have written a data race, "two potentially concurrent conflicting actions, at least one of which is not atomic at a scope that includes the thread that performed the other operation" (https://nvidia.github.io/cccl/unstable/libcudacxx/extended_api/memory_model.html , checked 2026-08-30).

One line apart, twice

Full program in code/day27-memory-model/memory_model.cu. It is the reduction above over 4,194,915 floats that are all 1.0f, run one hundred times per kernel. Every element being 1.0f is what makes the comparison exact rather than tolerant: every partial sum is a whole number under 2^24, where float loses nothing, and the signal to catch is one missing partial out of 1,024.

Two more rules apply.

The undefined kernels are reported and never gated. A data race is allowed to produce the right answer, so a check demanding the wrong one would assert that undefined behaviour is dependable. The gates are on the two defined kernels, exact on all one hundred runs.

Nothing is timed. No cudaEvent, no stopwatch, no number a clock produced. Day 42 is where a fence turns up as a stall reason in Nsight Compute.

Here is the hand-off with the fence left out:

    if (tid == 0) {
        partials[blockIdx.x] = blockSum;

        // The __threadfence() belongs on this line.

        const unsigned int ticket = atomicAdd(counter, 1u);
        isLastBlockDone = (ticket == gridDim.x - 1u);
    }
    __syncthreads();

And with it put back, which is the whole difference between the two rows of the first table:

    if (tid == 0) {
        partials[blockIdx.x] = blockSum;
        __threadfence();
        const unsigned int ticket = atomicAdd(counter, 1u);
        isLastBlockDone = (ticket == gridDim.x - 1u);
    }
    __syncthreads();

The barrier below the fence does a second job. Only thread 0 writes isLastBlockDone, so the rest of the block needs the barrier to read it, and because they then all read the same value the if below is uniform across the block, which is what makes the barriers inside the last block's second reduction legal. Day 14's rule.

There is a third version of those same lines that needs no separate fence, because the ordering rides on the atomic. It is libcu++, which the course reaches on day 39, and it is what NVIDIA's C++ language support page tells you to use.

cuda::atomic_ref<unsigned int, cuda::thread_scope_device> ticket{*counter};
partials[blockIdx.x] = blockSum;
// release publishes the store above; acquire hands every other block's
// store to whoever draws the last ticket.
const unsigned int old = ticket.fetch_add(1u, cuda::memory_order_acq_rel);
isLastBlockDone = (old == gridDim.x - 1u);

The second pair is one line each, on one address, from all 262,144 threads:

__global__ void tallyBlocksDeviceScope(unsigned int* tally) {
    atomicAdd(tally, 1u);
}

__global__ void tallyBlocksBlockScope(unsigned int* tally) {
    atomicAdd_block(tally, 1u);
}

Both are the contention case day 26 measures, and neither is timed here. The second is atomic inside a block and undefined across blocks, so the 256 threads of one block cannot lose an increment to each other and the 1,024 blocks can.

One honest caveat. The guide's worked sample marks the partial sums volatile on top of the fence; NVIDIA's shipping threadFenceReduction does not, and neither does this program. It relies on the fence alone, which is the choice the last pitfall below quotes the current guide to defend, and the hundred-run gate is what settles whether that choice held on this card.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5, 40 SMs)
n = 4194915 floats, every one of them 1.0f
1024 blocks x 256 threads, 100 launches per kernel
no timing anywhere in this program

reference total 4194915

kernel                runs  wrong  first bad run     total there
sumLastBlockRacy       100      0              -               -
sumLastBlockFenced     100      0              -               -

tally kernel              runs  wrong   lowest count        want
tallyBlocksDeviceScope     100      0         262144      262144
tallyBlocksBlockScope      100      0         262144      262144

reported, not gated: sumLastBlockRacy wrong on 0 of 100 runs, tallyBlocksBlockScope on 0 of 100

both gated kernels were exact on every run

The racy kernel produced the right answer on all 100 runs, but the memory model does not define it. The report keeps that result visible.

sumLastBlockRacy has no fence between writing its partial and announcing that the partial is ready. On this card and input, the write became visible before the announcement in every run. The memory model does not promise that.

A different GPU, compiler version, or launch configuration can change the order. The last block could then read an old partial value.

A test that passes 100 times out of 100 has not shown the code is correct. It has shown only that the bug did not appear. The memory model, not repeated success, decides whether the code is valid.

The fence defines the order used by the correct version. The two tally kernels differ in the same way. Blocks on different SMs must see each other's increments, so this code needs device scope.

The block-scope version reached 262,144 on all 100 runs but still has a data race across blocks. That result does not make the program valid.

Run it yourself

A free Colab T4 or any card you own. The program is one file and fits inside Compiler Explorer's twenty second cap, but the two edits the exercise asks for want a shell and a rebuild, so a notebook or a local card is the real path. Build line from the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o memory_model memory_model.cu

atomicAdd_block needs compute capability 6.0 or higher, which the course's sm_75 floor covers. Nothing in the file uses an API newer than CUDA 12.x.

Exercise

Break the fixed kernel two ways. Move __threadfence() from above the atomicAdd to below it and run three times. Then put it back, raise kBlocks from 1,024 to 8,192, and run the racy kernel again.

Time: 25 to 40 minutes. Submit: the wrong count for each of the three configurations, and one sentence on why a fence below the atomic is not a weaker fix but the same bug.

Check: two checks use branches that return EXIT_FAILURE, not an assert, because NDEBUG removes an assert from Release builds. The fenced kernel must equal the reference on every run, and the device-scope tally must read 262,144. A failure names the kernel, the number of failed runs, and the value read.

The program prints the racy rows but does not use them as checks. A zero wrong count there is a result, not a pass. Harness contract at /reference/harness.

Hint 1

Thread 0 does two things the last block depends on, in an order that matters: it publishes a number, and it announces the number is ready. Name both lines, then say what forces the publish to happen first.

Hint 2

Read what the guide promises for a legacy atomic. It promises atomicity, at a scope. Now ask what it promises about ordering, and what is left between the store and the atomic once you move the only line that promised anything.

Solution

A fence below the atomic orders the atomic against whatever you write after it, which nobody is waiting for. The two accesses needing an order are the store to partials[blockIdx.x] and the atomicAdd another block reads as permission to load that cell. Expect a wrong count like the racy kernel's, because that is the racy kernel with an extra instruction.

Raising kBlocks widens the window rather than creating it: at 1,024 blocks on a 40-SM card most blocks retired long before the last ticket was drawn.

The rule that outlives the exercise: a fence is an edge between two named accesses, not a property of a kernel. If you cannot say which store it orders against which announcement, you added an instruction rather than a fix, and day 14's barrier is the same sentence one scope down.

Pitfalls

Your multi-block reduction is right on 8 blocks and wrong on 8,192. With few blocks, the last counter update may occur after every other block has finished, so the missing order may not affect the result. Grid size does not make the code correct.

Test the grid size you plan to launch. A GPU with more SMs may expose the race more often.

You put the fence after the atomic. It compiles and it reads like a fix. It orders your writes against your later writes, and the write that mattered is above it.

You reached for atomicAdd_block because it is cheaper. The suffix is the scope, not a hint. It is atomic at cuda::thread_scope_block, so two blocks hitting one address with it are a data race, and the symptom is an undercount rather than a crash.

You forgot to reset the counter between launches. The second launch starts from a counter that already reached gridDim.x, no block draws gridDim.x - 1, and nothing writes the total. The guide's sample resets it inside the last block; this program zeroes it from the host, where you can watch it happen.

You reached for volatile. It is what most forum threads from before CUDA 9 tell you.

The current guide's own words: "The volatile keyword is supported to maintain compatibility with ISO C++. However, few, if any, of its remaining non-deprecated uses apply to GPUs", and "CUDA C++ volatile is NOT suitable for: Inter-Thread Synchronization: Use atomic operations via cuda::atomic_ref, cuda::atomic, or Atomic Functions instead." (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-support.html , section 5.3.10.4.3, checked 2026-08-30.)

Go deeper

Next

Day 28 replaces the ticket with a real grid-wide barrier through cudaLaunchKernelEx, and charges the price this pattern avoids: every block has to be resident, so a cooperative launch caps the grid below the one you just used. Day 29 takes the contention somewhere smaller instead of ordering it: a per-block copy of the counters in shared memory, where an atomic with 255 rivals is the fast path rather than the bug.