Triton #11334 — How Counter Values 0, 1, 2, and 3 Encode Two mbarrier Slots and Parity

Triton NVIDIA Barrier, Part 2 of 3 — Why reading the counter one step late can move a warp onto a different barrier generation

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.

Part 1 examined why Triton changed the ordering around the shared phase counter in NVIDIA’s cluster_barrier lowering.

Before the fix:

local barrier
→ counter load

After the fix:

counter load
→ local barrier

At the source level, the patch looked like a small reordering.

But the value being protected was not an ordinary loop counter.

The shared counter selected all of the following:

the physical mbarrier slot used by the current operation
the parity expected for that slot
the logical barrier generation that the warp should wait for

If two warps participating in the same source-level cluster_barrier read different counter values, they no longer wait on the same hardware-level barrier.

Two physical barrier slots represent four logical states

The relevant part of the lowering can be simplified to:

Value counter =
    load(counterPtr);

Value barrierIdx =
    counter & 1;

Value parity =
    counter >> 1;

Value nextCounter =
    (counter + 1) & 3;

Because nextCounter is masked with 3, the counter cycles through four values:

0 → 1 → 2 → 3 → 0

Those four values are interpreted as follows:

counter   binary   mbarrier slot   parity

0         00       slot 0          0
1         01       slot 1          0
2         10       slot 0          1
3         11       slot 1          1

The low bit selects the physical barrier slot:

counter & 1

The high bit selects the expected parity:

counter >> 1

A two-bit value therefore encodes two separate pieces of state:

Which physical mbarrier object should be used?

Which logical generation of that object should be awaited?

The counter is not merely a call number

It is easy to interpret counter values 0 and 1 as little more than:

first call
second call

But changing the counter from 0 to 1 changes the physical barrier address.

counter = 0
→ slot 0
→ parity 0

counter = 1
→ slot 1
→ parity 0

Changing the counter from 0 to 2 keeps the physical slot but changes the generation:

counter = 0
→ slot 0 / parity 0

counter = 2
→ slot 0 / parity 1

The two values refer to the same barrier storage, but not to the same barrier event.

The same barrier address does not necessarily mean the same barrier generation.

Why alternate between two physical slots?

Imagine using only one mbarrier slot for every cluster synchronization.

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 barrier while a slower CTA was still waiting for the previous round.

Slow CTA
→ still waiting for Round 0

Fast CTA
→ has already reused slot 0 for Round 1

The previous generation and the next generation would then occupy the same physical state too closely together.

Triton avoids immediate reuse by alternating between two slots:

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

A slot is reused only after the protocol has passed through the other slot.

slot 0
→ slot 1
→ slot 0

The source comment describes the same intent: a delayed CTA should not miss a phase because a peer reused one mbarrier twice before the delayed CTA began waiting.

Two slots alone are still not enough

Slot 0 is reused after two rounds:

Round 0 → slot 0
Round 2 → slot 0

The physical address is the same in both rounds.

Parity distinguishes those two logical uses:

Round 0
→ slot 0 / parity 0

Round 2
→ slot 0 / parity 1

Slot 1 follows the same pattern:

Round 1
→ slot 1 / parity 0

Round 3
→ slot 1 / parity 1

The protocol therefore forms a four-state ring:

slot 0 / parity 0
        ↓
slot 1 / parity 0
        ↓
slot 0 / parity 1
        ↓
slot 1 / parity 1
        ↓
slot 0 / parity 0

There are two physical barrier objects.

The combination of slot and parity creates four logical states.

Parity prevents an old completion from looking like a new one

Suppose slot 0 completed during an earlier round.

When the protocol later returns to slot 0, checking only:

Has slot 0 completed?

would be ambiguous.

The completion state from the previous use may still be visible.

Parity changes the question to:

Has slot 0 completed with the parity expected by this round?

These are different events:

slot 0 / parity 0 completed
≠
slot 0 / parity 1 completed

The wait operation therefore depends on both:

barrier pointer
+
expected parity

Two warps can calculate the same barrierPtr while still waiting for different generations if their parity values differ.

Reading the counter one step late always selects another slot

The low bit changes every time the counter advances by one.

Expanded into slot and parity:

0 → 1

slot 0 / parity 0
→
slot 1 / parity 0

1 → 2

slot 1 / parity 0
→
slot 0 / parity 1

2 → 3

slot 0 / parity 1
→
slot 1 / parity 1

3 → 0

slot 1 / parity 1
→
slot 0 / parity 0

A one-step difference does not merely move to the next numerical value.

It always selects the other physical slot.

On some transitions, parity changes as well.

What happens when one warp reads 0 and another reads 1?

Assume the current shared counter is 0:

counter = 0
→ slot 0
→ parity 0

One warp runs immediately after the old entry barrier.

The remaining warps are delayed before their counter loads.

Fast warp

loads counter 0
→ selects slot 0
→ expects parity 0

Delayed warps

have not loaded the counter yet

The fast path begins the current cluster-barrier protocol.

Assume the relevant peer CTAs also participate using:

slot 0 / parity 0

Once the current operation completes, thread 0 calculates:

nextCounter
=
(0 + 1) & 3
=
1

and stores it into shared memory:

shared counter:
0 → 1

A delayed warp now executes its first counter load:

loads counter 1
→ selects slot 1
→ expects parity 0

The same source-level barrier invocation now contains two views:

Fast warp
→ slot 0 / parity 0

Delayed warp
→ slot 1 / parity 0

They are no longer waiting for the same barrier generation.

The same function call can lower to different hardware waits

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

ttgl.barrier(cluster=True)

But the generated behavior can diverge:

Warp 0:
wait on mbarrier slot 0
with parity 0

Warp 1:
wait on mbarrier slot 1
with parity 0

The source-level operation is identical.

The hardware-level target is not.

Therefore:

reaching the same cluster_barrier call

does not by itself guarantee:

waiting on the same barrier generation

That guarantee exists only when all local participants derive their slot and parity from the same counter snapshot.

The wrong slot may receive no arrival for the current round

Suppose the peer CTAs send their current-round arrivals to:

slot 0 / parity 0

A delayed warp that selects slot 1 now waits on a barrier state that may receive no arrival for the current round.

Current peer arrivals
→ slot 0

Delayed warp’s wait
→ slot 1

Possible outcomes depend on timing and later activity:

the delayed warp may wait indefinitely
it may observe a future round’s signal as though it belonged to the current round
some warps may reach the final local barrier while others remain blocked
the CTA or CUDA context may deadlock

The exact failure does not need to be identical on every execution.

The protocol has already failed once current-round and next-round state are no longer separated.

The final local barrier cannot repair an earlier phase split

The old lowering also contained a local barrier after storing the next counter:

cluster wait
→ nextCounter store
→ tail local barrier

Could that final barrier reunify the warps?

No.

A warp that selected the wrong slot or parity may never finish its mbarrier wait.

Fast warp
→ correct wait completes
→ stores or observes the next state
→ reaches tail barrier

Delayed warp
→ waits on wrong slot or parity
→ does not complete
→ never reaches tail barrier

A local barrier releases only after all participating threads arrive.

If the barrier generation has already split before that point, the tail barrier cannot recover the lost agreement.

It may instead become another deadlock point.

A two-step difference selects the same slot but the opposite parity

Now suppose the expected counter is 0, but a thread reads 2.

Expected:
counter 0
→ slot 0 / parity 0

Observed:
counter 2
→ slot 0 / parity 1

The physical barrier pointer is the same:

slot 0

The expected generation is different:

parity 0
versus
parity 1

This demonstrates another important rule:

Calculating the same barrierPtr does not prove that two threads are waiting for the same barrier.

The logical generation is encoded separately.

One step late means the wrong slot; two steps late means the wrong parity

Counter drift can be summarized as follows.

One-step difference

counter + 1
→ another physical slot
→ sometimes a different parity as well

Two-step difference

counter + 2
→ the same physical slot
→ the opposite parity

Three-step difference

counter + 3
→ another slot
→ another logical generation

Any nonzero difference changes at least one part of the barrier identity.

same counter
→ same slot and parity

different counter
→ different slot, parity, or both

That is why coherent counter snapshots are a prerequisite for the entire protocol.

Double buffering and parity do not create snapshot consistency

The two-slot and parity scheme protects barrier reuse.

It does not automatically ensure that every local thread selects the same state.

The responsibilities are different.

Two physical slots
→ avoid immediately reusing the same barrier object in consecutive rounds

Parity
→ distinguish separate generations that reuse one slot

Coherent local counter snapshot
→ ensure all threads in the current call select the same slot and parity

The first two mechanisms cannot replace the third.

Double buffering is safe only when every participant agrees on which buffer is current.

The old ordering aligned the starting point, not the read completion

Before the fix:

local barrier
→ counter load

The local barrier guaranteed that all threads reached the instruction boundary before the load.

After the barrier released, scheduling could diverge again:

Warp 0
→ immediately executes the load

Warp 1
→ delayed by scheduling

Warp 2
→ waiting on another dependency

Warp 3
→ not selected yet

The old ordering proved:

Every participating thread reached the region before the counter load.

The protocol actually required:

Every participating thread completed the current counter load.

Those are different invariants.

The fixed ordering closes the reader set

After the fix:

counter load
→ local barrier

Each local thread first reads the published counter into a register.

Warp 0 → counter loaded
Warp 1 → counter loaded
Warp 2 → counter loaded
Warp 3 → counter loaded

Only then can all of them reach and release the local barrier.

all counter loads complete
        ↓
local barrier releases

Even if the shared counter changes later, every warp already owns the current snapshot in a private register.

shared counter:
0 → later becomes 1

current local snapshots:
all remain 0

After the barrier, every warp derives:

slot 0 / parity 0

from the same state.

A single writer does not guarantee coherent readers

Only thread 0 stores the next counter.

The lowering creates a leader predicate equivalent to:

threadId == 0

The next state is therefore written by one thread:

thread 0
→ stores nextCounter

all other threads
→ do not write it

This avoids multiple writers racing with one another.

But it does not automatically make all reads coherent.

single writer
≠
all readers observed the same version

Even with one writer, some readers can observe the old value while others observe the new value unless the read and write phases are ordered correctly.

The race in #11334 occurred at precisely that boundary.

The repaired protocol as a state transition

Assume the current counter is 0.

1. Publish the current state

shared counter = 0

The warp-specialization entry or previous tail barrier establishes this value for the new round.

2. Every local participant takes a snapshot

Warp 0 → reads 0
Warp 1 → reads 0
Warp 2 → reads 0
Warp 3 → reads 0

3. Seal completion of the reads

local barrier
→ no thread may continue until all snapshots exist

4. Select one barrier generation

Every warp derives:

slot 0
parity 0

5. Execute the cluster barrier

send or receive peer arrivals
+
wait for slot 0 / parity 0

6. Advance the shared state

Thread 0 stores:

nextCounter = 1

7. Publish the next state

The tail local barrier ensures that the next operation begins only after the current transition has completed locally.

The full sequence is:

publish
→ snapshot
→ seal
→ use
→ advance
→ publish

The counter is a compressed state machine

The shared integer represents a small state machine:

State 0
slot 0 / parity 0
        ↓
State 1
slot 1 / parity 0
        ↓
State 2
slot 0 / parity 1
        ↓
State 3
slot 1 / parity 1
        ↓
State 0

Every local thread participating in one transition must observe the same starting state.

If one thread observes State 0 while another observes State 1, one logical transition splits into two different actions:

same source operation
        ↓
different state snapshots
        ↓
different barrier selections

That is not a minor timing variation.

It is a divergence in the protocol’s state machine.

A small numerical difference can represent a large semantic difference

In ordinary arithmetic, the difference between 0 and 1 is small.

In encoded state, it can be structural.

0
→ physical barrier address A

1
→ physical barrier address B

Similarly:

0
→ one generation of slot 0

2
→ another generation of slot 0

The numerical difference may be one or two.

The execution difference is an entirely different wait condition.

This pattern appears often in low-level systems:

small integer
+
state-selection semantics
→ large behavioral difference

Reading the counter is therefore not just reading shared data.

It selects the identity of the current barrier operation.

What does it mean to wait on the same barrier?

The question cannot be answered only by checking whether threads reached the same source line.

They must agree on all of the following:

the same counter snapshot

the same physical mbarrier slot

the same expected parity

The actual contract is:

Threads participate in the same barrier only when they reach the same operation and derive the same physical slot and logical generation from the same published state.

If any one of those values differs, they are not waiting for the same barrier in execution terms.

Part 3: forcing the race instead of waiting for it

Under ordinary scheduling, every warp may read the counter quickly enough that the defect remains hidden.

The new regression test did not repeatedly launch the kernel and wait for a bad schedule to occur by chance.

It obtained the generated PTX and inserted a deliberate delay immediately before the counter load:

first warp
→ proceeds immediately

remaining warps
→ spin before the counter load

The test then compared the two orderings under the same forced schedule.

Old:
barrier → delayed load
→ fast warp can advance the counter

Fixed:
delayed load → barrier
→ fast warp must wait for every reader

Because the old behavior can deadlock the CUDA context, the dangerous launch is executed inside a timeout-bounded subprocess.

Part 3 examines how that PTX-level perturbation turned an intermittent scheduler race into a deterministic regression test.

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

Link the title above to Part 1 after publication.

Related material

Patch status: Merged into Triton main
Physical barrier slots: 2
Logical counter states: 4
State mapping: Low bit selects the slot; high bit selects parity
Core invariant: Every local participant in one cluster_barrier invocation must capture the same counter before thread 0 advances it

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

Part 1 explained why the old barrier → load ordering failed to prove that every warp had completed its phase read.

This article examined how four counter states encode two physical mbarrier slots and parity, and why reading the counter one step late can move a warp onto another barrier generation.

Part 3 examines how Triton modified generated PTX to delay selected warps at the exact counter-load boundary and isolated the potentially deadlocking launch inside a timeout-bounded subprocess.

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