Triton #11334 — Why cluster_barrier Could Advance the Phase Before Every Warp Had Read It

Triton NVIDIA Barrier, Part 1 of 3 — How moving a local barrier from before the load to after it restored the shared phase-snapshot contract

Update — August 18, 2026
The PTX delayed-counter dynamic regression test introduced with PR #11334 was removed in follow-up PR #11336 after timing out on GB200. The cluster_barrier compiler-ordering fix and the LLVM static regression checks remain in Triton main. References below to the dynamic test describe the test design that originally landed with #11334, not the test suite currently retained.

On August 17, 2026, Triton merged a fix for a phase race in the NVIDIA lowering of cluster_barrier.

The problem was not the complete absence of synchronization.

The old implementation already contained local barriers.

The problem was where one of those barriers was placed.

Before the fix, the relevant sequence was:

local barrier
→ load the shared phase counter
→ perform the cluster barrier
→ store the next counter
→ local barrier

At first glance, this appears properly synchronized.

Every thread reaches the first barrier before reading the counter, and every thread meets again at the final barrier.

But once the first barrier releases, there is no guarantee that every warp executes the following load at the same time.

One warp may read the counter immediately and continue through the barrier protocol.

Another warp may be delayed before it has read the counter at all.

If a fast thread advances far enough to store the next counter value, the delayed warp can read a different phase from the one used by the rest of the CTA.

The pull request summarized the race directly:

A thread could run ahead and overwrite the phase before another warp had finished reading it.

The fix reversed the ordering around the first local barrier:

load the shared phase counter
→ local barrier

That small change restored a much larger invariant:

Every participating local thread must capture the current phase before any thread is allowed to advance it.

What does `cluster_barrier` synchronize?

A normal CUDA thread block, or CTA, contains several warps.

Threads within that CTA can synchronize using a block-level barrier:

warps inside one CTA
→ local barrier

On NVIDIA Hopper and later architectures, several CTAs can also be grouped into a cluster.

CTA 0
CTA 1
CTA 2
CTA 3
→ one CTA cluster

A local CTA barrier cannot synchronize the entire cluster.

The participating CTAs need a wider protocol that communicates arrival and completion across CTA boundaries.

Triton’s NVIDIA lowering implements this path using mbarrier-based synchronization.

Conceptually:

each CTA reaches the cluster barrier
        ↓
arrival is communicated to peer CTAs
        ↓
the CTA waits for the current barrier phase
        ↓
all required CTAs arrive
        ↓
execution continues

But cluster-level synchronization is only one half of the problem.

The local warps inside each CTA must also agree on which barrier generation they are currently executing.

If different local warps read different phase state, they may wait on different physical barrier slots or different logical generations of the same slot.

Warp specialization creates a natural scheduling gap

The affected path is connected to warp specialization.

Under warp specialization, all warps in a CTA do not necessarily perform the same role.

For example:

Warp 0
→ producer or control role

Warp 1
→ worker role

Warp 2
→ worker role

Warp 3
→ worker role

Some warps may move data.

Others may perform computation.

Different partitions can have different dependencies and different execution timing.

Even after every warp has passed a local barrier, the scheduler does not promise to execute the next instruction of every warp simultaneously.

A possible schedule is:

Warp 0
→ local barrier releases
→ immediately loads the counter
→ continues through the cluster-barrier protocol

Warp 1
→ local barrier releases
→ delayed before loading the counter

Warp 2
→ local barrier releases
→ delayed before loading the counter

Warp 3
→ local barrier releases
→ delayed before loading the counter

The local barrier proves that every participating thread reached the same point.

It does not prove that every thread completed the instruction after that point.

That distinction is the center of the race.

The counter encodes both barrier slot and parity

The lowering stores a small counter in shared memory.

In simplified form, the code derives the barrier state like this:

Value counter = load(counterPtr);

Value barrierIdx =
    counter & 1;

Value parity =
    counter >> 1;

After the barrier operation completes, the counter advances:

Value nextCounter =
    (counter + 1) & 3;

The & 3 keeps the counter within four states:

0 → 1 → 2 → 3 → 0

Those four states encode two separate pieces of information.

low bit
→ physical mbarrier slot

high bit
→ logical parity of that slot

Expanded:

counter = 0
→ slot 0
→ parity 0

counter = 1
→ slot 1
→ parity 0

counter = 2
→ slot 0
→ parity 1

counter = 3
→ slot 1
→ parity 1

The counter is therefore not merely a call count.

It identifies the physical barrier object and the logical generation that the current call must wait for.

Why alternate between two barrier slots?

Imagine reusing only one mbarrier slot for every cluster barrier.

Round 0
→ slot 0

Round 1
→ slot 0 again

Round 2
→ slot 0 again

A fast CTA could complete one round and begin reusing the same physical barrier while another CTA is still waiting on the previous round.

slow CTA
→ still waiting for Round 0

fast CTA
→ has already reused slot 0 for Round 1

The old and new generations could become difficult to distinguish.

Alternating between two slots creates space between consecutive uses of the same physical object:

Round 0 → slot 0
Round 1 → slot 1
Round 2 → slot 0
Round 3 → slot 1

Parity then distinguishes separate generations that reuse the same slot:

Round 0
→ slot 0 / parity 0

Round 2
→ slot 0 / parity 1

This double-buffered protocol is useful only if every local participant agrees on the same counter.

If one warp reads 0 and another reads 1, they do not merely disagree about a number.

They select different physical barrier slots.

The old implementation read the counter after the local barrier

Before the patch, the lowering contained the equivalent of:

NVVM::BarrierOp::create(
    rewriter,
    loc
);

Value counter =
    b.load(
        i32_ty,
        counterPtr
    );

The execution order was:

entry local barrier
        ↓
each thread loads the counter
        ↓
slot and parity are calculated
        ↓
cluster mbarrier arrival and wait
        ↓
the leader stores nextCounter
        ↓
tail local barrier

The first barrier placed all local threads at the same starting point.

But it did not seal the counter read.

After the barrier released, one warp could execute the shared-memory load while another remained unscheduled.

The protected operation was on the wrong side of the synchronization point.

A barrier closes what comes before it

A simplified barrier guarantee is:

No participating thread may continue beyond this point until every participating thread has reached it.

The placement of the barrier determines what that guarantee covers.

operations before the barrier
→ completed by every participant before release

operations after the barrier
→ not yet guaranteed to have completed

The old ordering:

barrier
→ counter load

proved:

Every local thread reached the point immediately before the counter load.

But the protocol required a stronger statement:

Every local thread has already loaded the current counter into its own register.

Those are not equivalent guarantees.

A possible race, step by step

Assume the shared counter currently contains 0.

counter = 0
→ slot 0
→ parity 0

1. Every warp reaches the entry barrier

Warp 0 → arrived
Warp 1 → arrived
Warp 2 → arrived
Warp 3 → arrived

The local barrier releases them.

2. Warp 0 runs ahead

Warp 0
→ loads counter 0
→ selects slot 0 / parity 0
→ begins the cluster-barrier protocol

The remaining warps have passed the entry barrier but have not yet executed the load.

Warp 1 → counter not loaded
Warp 2 → counter not loaded
Warp 3 → counter not loaded

3. The fast path completes the current phase

If the required peer arrivals are already available, the fast thread can complete the cluster wait.

The leader then computes:

nextCounter
=
(0 + 1) & 3
=
1

and stores it into shared memory:

shared counter:
0 → 1

4. A delayed warp finally executes the load

Warp 1
→ loads counter 1
→ selects slot 1 / parity 0

Now one logical cluster_barrier invocation has split into two local views:

Warp 0
→ slot 0 / parity 0

Warp 1
→ slot 1 / parity 0

They have reached the same source-level barrier operation.

They are no longer waiting on the same hardware-level barrier generation.

Why the final local barrier cannot repair the split

The old code also contained a local barrier after the leader stored the next counter.

Could that final barrier bring the warps back into agreement?

No.

Once a delayed warp selects the wrong slot or parity, it may block inside a different mbarrier wait.

fast warp
→ waits on slot 0 / parity 0
→ completes
→ reaches the tail local barrier

delayed warp
→ waits on slot 1 / parity 0
→ current round may never complete that slot
→ never reaches the tail local barrier

A local barrier can release only when all required participants reach it.

If some warps are already trapped in a different barrier generation, the final barrier is too late to restore the protocol.

It can instead become another point of deadlock.

The fix moved the barrier after the load

After the patch, the relevant order became:

Value counter =
    b.load(
        i32_ty,
        counterPtr
    );

NVVM::BarrierOp::create(
    rewriter,
    loc
);

Or:

load the shared counter
→ local CTA barrier

The source diff is small.

The guarantee changes completely.

The new ordering means:

every thread loads the current counter
        ↓
every thread reaches the local barrier
        ↓
only then may any thread continue

Once the barrier releases, every local thread already owns a private register snapshot of the same published counter.

The shared value can change later without changing the current call’s local snapshots.

The new order establishes read-before-advance

The corrected protocol can be written as:

the current counter is published
        ↓
every local thread reads it
        ↓
a local barrier seals completion of all reads
        ↓
every thread derives the same slot and parity
        ↓
the cluster barrier executes
        ↓
the leader stores the next counter
        ↓
the tail local barrier publishes the next state

Or more compactly:

publish
→ read
→ rendezvous
→ use
→ advance
→ publish

The central invariant is:

The writer may advance the shared state only after every reader has captured the current state.

Why is another barrier not required before the load?

Moving the barrier after the load raises a natural question:

What prevents a thread from reading while the counter is still being updated from the previous round?

The implementation’s comment identifies the publication boundary.

The current counter has already been published by either:

the warp-specialization entry synchronization

or:

the previous cluster barrier’s tail local barrier

The previous round ends with:

leader stores nextCounter
→ tail local barrier

That tail barrier ensures the next round does not begin while the counter update is still incomplete.

The current round can therefore begin by loading the published counter.

The following local barrier serves a different purpose: it ensures that every local reader has completed that load before the state can advance again.

previous tail barrier
→ publishes the write

current post-load barrier
→ seals the reads

The two barriers protect different boundaries.

The patch did not add synchronization. It corrected its placement.

It would be inaccurate to summarize this bug as:

The implementation forgot a barrier.

The old implementation already had two local barriers.

The real problem was:

A barrier existed,
but it was placed on the wrong side of the operation it needed to protect.

Compare:

barrier → load

with:

load → barrier

The first ordering aligns the point at which threads are allowed to begin the load.

The second ordering aligns the point at which all threads have completed the load.

Because the shared counter could be advanced afterward, the protocol required the second guarantee.

Why was the race difficult to observe during normal execution?

Under ordinary scheduling, all warps may load the shared counter almost immediately after the entry barrier releases.

A shared-memory load is short, and the fast warp may not have enough time to complete the cluster operation and update the counter before the other warps read it.

typical scheduling
→ every warp reads the counter quickly
→ every warp happens to observe the same phase
→ test passes

But concurrency correctness cannot depend on the scheduler usually choosing a convenient order.

The relevant questions are:

Can a warp be delayed before the load?
→ Yes.

Can another thread advance the counter while it is delayed?
→ Under the old ordering, yes.

A race does not have to occur on every execution.

It is enough that one valid execution order can violate the shared-state contract.

The regression test forces the dangerous schedule

The new test does not repeatedly run the kernel and hope that the GPU scheduler eventually creates the race.

It first obtains the generated PTX.

It then locates the exact shared-counter load and deliberately delays every warp except the first immediately before that instruction.

Conceptually:

first warp
→ reaches the counter load immediately

remaining warps
→ delayed immediately before the counter load

This is more precise than inserting a delay in the source code.

A source-level delay placed before cluster_barrier() would occur before the lowering-generated entry barrier.

The first warp would simply wait at that entry barrier for the delayed warps, preventing the desired race window.

The test therefore modifies the generated PTX at the exact instruction boundary where the old protocol became unsafe.

Part 3 of this series examines that regression test in detail.

The test is isolated because the failure can deadlock the CUDA context

This race may not produce a clean numerical mismatch.

If different warps wait on different barrier slots or phases, some of them may never complete.

The CUDA context can become stuck.

A normal assertion cannot report a result if the kernel call never returns.

The regression test therefore runs the deliberately delayed kernel in a separate process with a bounded timeout.

parent test process
        ↓
isolated child process
        ↓
delayed GPU kernel
        ↓
normal completion or timeout

The child can be terminated without leaving the entire test suite permanently blocked.

This is another important distinction:

incorrect value
→ ordinary assertion may be enough

possible nontermination
→ process isolation and timeout are part of the test contract

One reordered instruction restored a phase-snapshot invariant

At the source-diff level, the fix moved a barrier across a load.

Before

barrier();
counter = load();

After

counter = load();
barrier();

But the repaired invariant is larger than the diff:

Every local thread participating in one cluster_barrier invocation must capture the same counter value before the leader may advance that value.

Without this invariant, the two-slot and parity protocol cannot remain coherent.

Adding more slots would not solve the underlying race.

Using a larger counter would not solve it either.

If readers do not agree on the current state, the same source-level call can still split into different hardware generations.

A barrier contract includes its position

Seeing a barrier in low-level code does not prove that the protected state is safe.

The review must also ask:

Which write publishes the state?

Which load consumes it?

Is the barrier before or after that load?

Which threads participate?

Can one participant advance the state while another has not observed it?

In this case, the number of barriers was not the central issue.

Their relationship to the shared counter was.

The necessary condition was not merely:

Has every thread reached this code region?

It was:

Has every thread captured the current phase?

Reaching the same function does not mean waiting on the same barrier

At the source level, every warp executes the same operation:

cluster_barrier

But if their counter snapshots differ, the resulting waits differ:

Warp 0
→ slot 0 / parity 0

Warp 1
→ slot 1 / parity 0

The function call is the same.

The physical slot or logical generation is not.

The central lesson is therefore:

Participating in the same barrier means more than reaching the same instruction. It means using the same phase snapshot.

Related material

Patch status: Merged into Triton main
Merged: August 17, 2026 UTC
Affected path: NVIDIA cluster_barrier lowering under warp specialization
Regression-test hardware scope: NVIDIA Hopper or later
Possible failure mode: Local warps observe different phase state, split across barrier generations, and may deadlock
Restored invariant: Every participating local thread must read the current counter before the leader can store the next one

This is Part 1 of a three-part series on Triton’s NVIDIA cluster_barrier phase race.

Part 2 examines how counter values 0, 1, 2, and 3 encode two physical mbarrier slots and two parity generations — and why reading the counter only one step late can move a warp onto an entirely different barrier.

#Triton #NVIDIA #CUDA #Hopper #WarpSpecialization #ClusterBarrier #GPUConcurrency #RaceCondition #CodeAnalysis #SoftwareArchitecture

Search This Blog

Resonetic Lab | AI · Code · Market