Day 75Module 8

draft

Asynchronous barriers

This page uses four warps to load tiles and four warps to compute. The loader warps finish each copy in a few hundred cycles, then wait at __syncthreads() while the other warps run 64 chained FMAs per element.

__syncthreads() makes all eight warps wait for the slowest warp. Day 39 introduced libcu++, which exposes the hardware alternative as cuda::barrier, available from compute capability 8.0. This lesson splits the barrier into arrive and wait operations, then uses them in a producer-consumer schedule.

One barrier, two separable promises

__syncthreads() combines two guarantees: no thread continues until every thread arrives, and every write made before the call is visible after it. A thread with independent work must still wait.

An asynchronous barrier separates those operations. token = bar.arrive() announces this thread's part is done, and bar.wait(std::move(token)) blocks until everyone has announced. The CUDA 12.6 programming guide, section 7.26.2: "the call to bar.arrive() does not block a thread, it can proceed with other work that does not depend upon memory updates that happen before other participating threads' call to bar.arrive()".

The memory guarantee applies to the pair: updates made before an arrive are guaranteed visible after the matching wait (https://docs.nvidia.com/cuda/archive/12.6.2/cuda-c-programming-guide/index.html#asynchronous-barrier , checked 2026-09-01). Between its arrive and its wait, a thread is free.

The object behind it lives in shared memory and, from compute capability 8.0, it is a hardware unit. The PTX ISA calls it an mbarrier: "a barrier created in shared memory", one 8-byte .b64 word tracking the phase and the count of pending arrivals, and "mbarrier operations enable threads to perform useful work after the arrival at the mbarrier and before waiting for the mbarrier to complete". mbarrier.init, mbarrier.arrive and mbarrier.test_wait each carry the Target ISA note "Requires sm_80 or higher", and the transaction-counting variants day 76 needs are sm_90 (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier , checked 2026-09-01).

Below that the same C++ compiles and runs, but in software: "Devices of compute capability 8.0 or higher provide hardware acceleration for barrier operations and integration of these barriers with the memcpy_async feature", while "On devices with compute capability below 8.0 but starting 7.0, these barriers are available without hardware acceleration" (guide section 7.26, at the 12.6 guide link above). This lesson covers the hardware implementation, not the software fallback.

Diagram: one tile handoff, two schedules. Two horizontal bands, each showing 8 warp timelines in one block: 4 producer warps above 4 consumer warps, time flowing right, with tile fill bars on producers and long compute bars on consumers. Band 1, __syncthreads: all 8 timelines stop at a vertical line twice per tile; producers show dead space from fill-end to the line. Caption "one all-stop per boundary: 8 warps wait for the slowest". Band 2, split barrier: producers hit a dot (arrive) at fill-end and keep going to the next ready wait; consumers hit their own dot after draining. Caption "one-sided waits: each group blocks only on the promise it needs, 4 barrier objects, 2 buffers". Alt text: "Two schedules of four producer and four consumer warps. Syncthreads stops all eight warps at every tile boundary. The split barrier lets a warp arrive, keep working, and wait only when it next needs the other group."

The arrive call in code, from the guide's own pattern, is one line each way:

barrier::arrival_token token = bar.arrive();  // does not block
independentWork();
bar.wait(std::move(token));  // blocks until all have arrived

A barrier is an object, not a place

Every barrier used so far was a program point that all threads had to reach together. __syncthreads requires this rule, which is why day 14 banned it inside a divergent branch and day 65 showed the resulting hang.

Day 14 uses __syncthreads as the block-wide baseline for this comparison.

cuda::barrier breaks the model on purpose. It is an object in shared memory, and arriving at it is just an operation on that object, legal from any thread at any line. The guide's producer-consumer pattern (section 7.26.5, "Spatial Partitioning (also known as Warp Specialization)") puts the two thread groups in two different branches of one if, each looping over its own arrive and wait calls, a shape that would deadlock instantly with __syncthreads.

The waits are one-sided by design: "Producer threads wait for consumer threads to signal that the buffer is ready to be filled; however, consumer threads do not wait for this signal. Consumer threads wait for producer threads to signal that the buffer is filled; however, producer threads do not wait for this signal" (same guide section, checked 2026-09-01).

These one-sided waits implement warp specialization. The guide also gives the sizing rule used here: "For full producer/consumer concurrency this pattern has (at least) double buffering where each buffer requires two cuda::barriers". Two buffers, four barriers, day 54's trade relocated inside one block.

Two kernels, one schedule

Full program in code/day75-async-barriers/barrier_pipeline.cu. Both kernels walk the same tiles with the same grid, the same 4-and-4 warp split and the same double-buffered shared array, so what holds the two warp groups together is the only variable. Correctness comes before any clock: the __syncthreads version is checked against a CPU reference, the output is zeroed, and the split version must then memcmp equal, so a kernel that skipped its stores cannot look fast.

The warp split itself is the constant kProducerWarps at the top of the file, because the exercise moves it.

// The same schedule on four cuda::barrier objects, the producer-consumer
// pattern of programming guide section 7.26.5. ready[b] means "tiles[b]
// may be overwritten"; filled[b] means "tiles[b] holds a complete tile".
// Each wait is one-sided: producers never wait on filled, consumers never
// wait on ready. arrive() returns a token that is the right to wait on
// that phase; a group that will never wait drops it, and the (void) says
// the drop is deliberate.
//
// All 256 threads participate in all four barriers, which is why each is
// initialized with kThreadsPerBlock: an arrive_and_wait and a bare arrive
// count the same. On sm_80 and newer these barriers are hardware mbarrier
// objects in shared memory.
__global__ void pipelineSplitBarrier(const float* __restrict__ in,
                                     float* __restrict__ out, size_t n) {
    __shared__ float tiles[2][kTileElems];
    __shared__ barrier ready[2];
    __shared__ barrier filled[2];

    if (threadIdx.x < 2) {
        init(&ready[threadIdx.x], kThreadsPerBlock);
        init(&filled[threadIdx.x], kThreadsPerBlock);
    }
    __syncthreads();  // no thread may touch a barrier before init

    const size_t numTiles = (n + kTileElems - 1) / kTileElems;

    if (threadIdx.x < kProducerThreads) {
        size_t j = 0;
        for (size_t tile = blockIdx.x; tile < numTiles;
             tile += gridDim.x, ++j) {
            const int buf = static_cast<int>(j & 1);
            ready[buf].arrive_and_wait();  // consumers done with tiles[buf]
            loadTile(in, tiles[buf], tile, n);
            (void)filled[buf].arrive();  // hand off; do not wait
        }
    } else {
        (void)ready[0].arrive();  // both buffers start out writable
        (void)ready[1].arrive();
        size_t j = 0;
        for (size_t tile = blockIdx.x; tile < numTiles;
             tile += gridDim.x, ++j) {
            const int buf = static_cast<int>(j & 1);
            filled[buf].arrive_and_wait();  // producers filled tiles[buf]
            consumeTile(tiles[buf], out, tile, n);
            (void)ready[buf].arrive();  // hand back; do not wait
        }
    }
}

Staging a plain array through shared memory does not improve a copy. The workload isolates the schedule: replace the FMA chain with a fragment multiply and loadTile with cp.async, and it becomes day 74's pipelined GEMM. The fixed cost of 64 dependent FMAs per element keeps the schedules comparable.

Results

Not yet run. This binary targets sm_80 and its capability gate skips the software fallback. Run it on a CC 8.0 or newer GPU and save the transcript in code/day75-async-barriers/evidence/. Until then, every hardware field in the front matter stays null.

Kernel	Mean (ms)	split/sync
pipelineSyncthreads	?
pipelineSplitBarrier	?	?

The first run must check three claims:

Both correctness gates pass. The CPU check and the memcmp are the harness. A mismatch means a barrier is in the wrong place.
split/sync lands at or below about 1.0. The gap should be modest because the baseline already double-buffers, so fill-compute overlap exists in both kernels. What the split removes is the twice-per-tile all-stop and the requirement that both groups drain to the boundary in lockstep. If the ratio comes back clearly above 1.0, the two extra barrier operations per group per tile cost more than the lockstep did, and this page will say so in this table.
The direction survives a re-split. The exercise's 2-producer-warp build must move both kernels the same direction, or the explanation below is wrong.

Compare the ratio and schedule, not the absolute times. SM counts and clocks differ across CC 8.0 cards; whether one-sided waits beat all-stops at this arithmetic intensity is the part that transfers.

Run it yourself

This lesson needs compute capability 8.0 or newer. Every GeForce card since the RTX 30 series qualifies (8.6, 8.9, 12.0), as does any A100 or H100. The README lists the build command and expected artifacts.

The setup guide covers remote GPU options.

Exercise

Rebuild with kProducerWarps = 2 (two loader warps, six math warps). Before running, write down which direction split/sync and both absolute times should move, and why. Then run and check yourself.

Time: 25 to 40 minutes. Submit: the two split/sync ratios (4-warp and 2-warp builds), both mean times, and one sentence naming the group that sets the pace.

Check: the program is the harness. On pass it prints both means and the ratio; on a broken schedule it prints either the first mismatching index against the CPU reference or split-barrier output differs from the __syncthreads baseline and exits nonzero. Both static_asserts hold at any whole-warp split, so a build that compiles is a legal configuration.

Hint 1

Per tile, one producer thread issues a handful of loads while one consumer thread runs a handful of 64-FMA chains. Which group finishes first today, and what does the parked group's idle time cost under each barrier scheme?

Hint 2

At 2-and-6 each producer thread now covers 16 elements of the tile and each consumer thread about 5. The question is not just which kernel wins, but which kernel loses less when the two groups' costs move closer together. Where does an all-stop hurt most: balanced groups or lopsided ones?

Solution

The consumer group sets the pace because its per-tile arithmetic costs more than the producer's loads. Shifting two warps from loading to math should make both kernels faster.

With a 4-and-4 split, __syncthreads stalls the faster group at every boundary, so the split barrier has more idle time to recover. At 2-and-6 the group costs are closer and the two kernels should converge. If the ratios order the other way, barrier overhead is the larger term on that GPU.

A split barrier reduces the wait time of the group that finishes early.

Pitfalls

A thread touched a barrier before it was initialized. init runs on one thread, and nothing orders other threads behind it unless you put a __syncthreads() after the init block. The PTX contract is blunt: "Performing any mbarrier operation except mbarrier.init on an uninitialized mbarrier object results in undefined behavior". No error code, just a broken phase count.

The producer writes the buffer after arriving. The fence travels with the pair: only writes made before the arrive are visible after the matching wait. A store between your arrive and the consumers' wait is a data race that passes on a quiet card. Day 62's racecheck is how you find it after the fact.

The init count does not match the arrivals. Each phase completes when the count reaches the expected arrival count passed to init, and an arrive_and_wait counts exactly one arrive. Initialize a block-participation barrier with 128 while 256 threads arrive and the phase flips halfway through, releasing waiters while the other group is mid-write. Count arrivals per phase, not threads you think of as participants.

You carried the split structure back to __syncthreads. The two-branch producer-consumer shape is legal for barrier objects and a guaranteed hang for __syncthreads, which every thread of the block must reach. That is day 65's first hang, reproduced in one refactor.

You measured cuda::barrier on a T4 and published the number. It runs: the barriers work in software from CC 7.0. But below 8.0 they run "without hardware acceleration" (guide section 7.26), so the number prices a software loop, not the mbarrier. This program gates on CC 8.0 so that mistake cannot ship; name the card next to any barrier cost you quote.

Go deeper

CUDA C++ Programming Guide 12.6, section 7.26, "Asynchronous Barrier": https://docs.nvidia.com/cuda/archive/12.6.2/cuda-c-programming-guide/index.html#asynchronous-barrier (checked 2026-09-01). The current guide carries it at https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html (checked 2026-09-01)
PTX ISA, "Parallel Synchronization and Communication Instructions: mbarrier": https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier (checked 2026-09-01)
libcu++ (CCCL) documentation, cuda::barrier and friends: https://nvidia.github.io/cccl/unstable/libcudacxx/ (checked 2026-09-01)

Day 76 uses a hardware copy engine as the producer: TMA unit copies whole tiles from a single thread's request, and it reports completion by adding a transaction count to this same mbarrier, using its byte-counting variant. Day 80 uses this pattern with day 74's cp.async as the producer and tensor core math as the consumer, timed with CUDA events.