Triton #11334 and #11336 — Why Was the PTX Delay Test for the cluster

Triton NVIDIA Barrier, Part 3 of 3 — A deterministic reproduction attempt, a GB200 timeout, and the regression contract that remains

The first two articles in this series examined a phase race in Triton’s NVIDIA cluster_barrier lowering.

The old implementation used this order:

local barrier
→ load the shared counter

The local barrier proved that every participating warp had reached the point before the counter load.

It did not prove that every warp had actually completed the load.

A fast warp could read the current counter, pass through the cluster-barrier protocol, and allow thread 0 to publish the next counter before a delayed warp had captured the old value.

One CTA could then split across two barrier generations:

Fast warp
→ counter 0
→ slot 0 / parity 0

Delayed warp
→ counter 1
→ slot 1 / parity 0

PR #11334 changed the order to:

load the shared counter
→ local barrier

A fast warp can still perform the load first.

But it cannot continue beyond the following local barrier until every other local warp has also captured its counter snapshot.

PR #11334 did more than repair the compiler ordering.

It also introduced a generated-PTX test designed to force the dangerous schedule deliberately.

The test found the exact counter-load instruction, allowed the first warp to proceed immediately, delayed the remaining warps immediately before that load, and ran the potentially deadlocking kernel inside a timeout-bounded subprocess.

That dynamic test did not remain in the repository for long.

Follow-up PR #11336 removed it after the test timed out on GB200. The Python test file was restored to its pre-#11334 form, while the compiler fix and LLVM instruction-order regression checks were explicitly retained.

This article is therefore not a description of a dynamic PTX-delay test currently retained in Triton main.

It examines:

how the removed test attempted to reproduce the race deterministically,
why a source-level delay was not precise enough,
what the GB200 timeout does and does not establish,
which regression evidence remains today, and
what coverage was lost when the dynamic test was removed.

An ordinary functional test could not pin down the race

The test kernel used two warp-specialized partitions.

Each partition crossed a cluster barrier and then wrote a different half of an output tensor.

Conceptually:

Partition 0
→ cluster barrier
→ write output[0:128]

Partition 1
→ cluster barrier
→ write output[128:256]

A successful launch should produce:

0, 1, 2, 3, ... , 255

The ordinary Hopper-or-later functional test launches the kernel with two or four CTAs in the cluster and checks that:

the expected cluster mbarrier instructions were generated

an unexpected mapa path was not generated

the kernel completed

the output matched 0 through 255

That is useful evidence.

It shows that the current cluster-barrier path works under an ordinary execution schedule.

But it does not directly answer the question behind the race:

Does the protocol remain correct when one local warp is delayed for a long time immediately before reading the shared counter?

A race condition does not have to fail on every execution.

Incorrect synchronization may still pass whenever the scheduler happens to produce a convenient order.

Incorrect synchronization
+
A safe-looking schedule
→ PASS

Repeatedly launching the ordinary kernel might eventually expose the race.

It might also run thousands of times without doing so.

The defect depended on a specific ordering between fast and delayed readers.

The regression test therefore needed to construct that ordering directly.

The schedule the test tried to create

Before the fix, the relevant lowering order was:

entry local barrier
→ shared counter load

Assume one CTA contains four warps:

Warp 0
Warp 1
Warp 2
Warp 3

Once all four warps reach the entry barrier, the barrier releases.

The scheduler is then free to execute the next instruction from one warp while delaying the others.

The dangerous schedule is:

Warp 0
→ immediately load the current counter
→ enter the cluster-barrier protocol
→ progress far enough for thread 0 to publish nextCounter

Warp 1
→ has not loaded the counter yet

Warp 2
→ has not loaded the counter yet

Warp 3
→ has not loaded the counter yet

If the shared counter begins at 0, Warp 0 selects:

counter 0
→ slot 0
→ parity 0

If thread 0 advances the shared state before the other warps perform their loads, those delayed warps can observe:

counter 1
→ slot 1
→ parity 0

Reproducing this condition requires more than making some warps slow.

The delay must occur at a precise boundary:

after the old entry barrier
+
immediately before the shared counter load

Placed slightly earlier, the entry barrier reunites the warps.

Placed slightly later, all warps may already have captured the same counter.

Why a source-level delay was not precise enough

The most obvious approach would be to add a delay to the Triton or Gluon kernel:

if not_first_warp:
    delay()

ttgl.barrier(cluster=True)

But the local barrier and counter load were generated internally by the compiler lowering of cluster_barrier().

A delay written before the source-level barrier operation could lower into:

source-level delay
→ compiler-generated entry barrier
→ counter load

Warp 0 could skip the delay and reach the generated entry barrier first.

But it would then wait there for Warps 1, 2, and 3.

Warp 0
→ no delay
→ reaches the entry barrier
→ waits for the delayed warps

Once the remaining warps finish their delays and reach the entry barrier, all of them are released into the counter-load region.

The state required for the race does not appear:

some warps have loaded the counter

some warps have not loaded it

the fast thread can already advance the shared state

The test did not merely need “a slow warp.”

It needed a slow warp inside the lowering-generated protocol, immediately before the counter load.

That is why the test moved below the source level and modified the generated PTX.

The test first captured the generated PTX

PR #11334 compiled the kernel through a warmup path and obtained the generated PTX string.

Conceptually:

original_ptx = kernel.warmup(...).asm["ptx"]

It then searched for the local barrier and shared counter load.

The old ordering appeared as:

bar.sync
ld.shared counter

The fixed ordering appeared as:

ld.shared counter
bar.sync

The regular expression was designed to recognize both forms:

Old:
barrier
load

or:

Fixed:
load
barrier

That design was intentional.

The same test structure was meant to reason about both sides of the patch:

Old ordering
→ create a window in which a fast warp may advance the counter

Fixed ordering
→ verify that the post-load barrier prevents that advance

There is one evidentiary boundary worth preserving.

The source confirms that the test was written to recognize old and fixed instruction sequences.

The public record does not, by itself, prove that a negative-control run was completed against the old compiler commit and reliably reproduced a timeout there.

The design intent is visible.

A completed old-versus-fixed experimental matrix is not fully documented in the material examined here.

The delay was always inserted immediately before the counter load

After finding the instruction pair, the test located the beginning of the shared counter load.

It inserted the PTX delay block directly before that instruction.

Under the old code:

barrier
→ injected delay
→ counter load

Under the fixed code:

injected delay
→ counter load
→ barrier

The semantic insertion point stayed the same:

Pause selected warps immediately before they read the current shared counter.

The difference was what the fast warp could do after its early load.

Only the first warp skipped the delay

Assume the CTA has four warps and 128 threads:

Warp 0
→ thread IDs 0–31

Warp 1
→ thread IDs 32–63

Warp 2
→ thread IDs 64–95

Warp 3
→ thread IDs 96–127

The injected PTX used %tid.x to separate the first warp from the others.

thread ID < 32
→ skip the delay

thread ID >= 32
→ enter the spin loop

The resulting schedule was:

Warp 0
→ fast reader

Warps 1–3
→ delayed readers

The delay used the GPU’s clock64 value.

Conceptually:

start = clock64

while clock64 - start < 10,000,000:
    continue

The number 10,000,000 should not be interpreted as one portable duration in seconds.

GPU clock behavior varies across architectures and runtime conditions.

The purpose was simply to create a large enough scheduling gap for Warp 0 to move far ahead of the other local warps.

The intended race window under the old ordering

Assume the shared counter begins at 0.

1. Every warp passes the entry barrier

Warp 0 → released
Warp 1 → released
Warp 2 → released
Warp 3 → released

2. The injected delay begins

Warp 0
→ skips the delay
→ loads counter 0

Warps 1–3
→ spin immediately before the counter load

3. Warp 0 selects the current generation

counter 0
→ slot 0
→ parity 0

If the peer CTAs also progress through that generation, the current cluster barrier can complete for the fast path.

4. Thread 0 advances the shared counter

nextCounter
=
(0 + 1) & 3
=
1

shared counter:
0 → 1

5. The delayed warps finish spinning

Warp 1 → loads counter 1
Warp 2 → loads counter 1
Warp 3 → loads counter 1

One CTA now contains two barrier identities:

Warp 0
→ slot 0 / parity 0

Warps 1–3
→ slot 1 / parity 0

Some warps can now wait on a barrier generation that the current round does not complete.

The result may be more severe than an incorrect numerical value.

It can become kernel nontermination or a stuck CUDA context.

The fixed ordering should stop the fast warp at the local barrier

After the fix, the sequence is:

counter load
→ local barrier

The test still inserts the delay immediately before the counter load.

1. Warp 0 reads the counter first

Warp 0
→ no delay
→ loads counter 0

2. Warp 0 reaches the post-load barrier

Warp 0
→ arrives at the local barrier
→ waits for the other warps

It cannot enter the cluster mbarrier protocol alone.

Thread 0 therefore cannot progress to the point where it publishes nextCounter.

3. The delayed warps finish spinning

Warp 1 → loads counter 0
Warp 2 → loads counter 0
Warp 3 → loads counter 0

They then reach the same local barrier.

4. The barrier releases only after all snapshots exist

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

Every warp selects:

slot 0 / parity 0

5. The shared counter advances only afterward

The cluster barrier completes.

Thread 0 eventually stores 1.

By then, every local warp already holds counter 0 in its register.

Changing the shared-memory value no longer changes the identity of the current barrier operation.

That is the read-before-advance invariant restored by:

load
→ barrier

The test also checked that the modified PTX was actually used

Creating a modified PTX string was not enough.

The test wrote the delayed PTX to a temporary file and invoked the kernel through Triton’s ir_override path.

Conceptually:

compiled = kernel[...](
    ...,
    ir_override=delayed_ptx,
)

It then compared the PTX associated with the compiled object against the modified PTX.

compiled.asm["ptx"] == delayed_ptx

Without this check, the following false PASS would be possible:

The test creates delayed PTX

A cache or compilation path ignores the override

The ordinary PTX executes

The kernel completes

The test incorrectly reports success

The test tried to prove not only that the perturbation had been constructed, but that the perturbed code was the code selected for execution.

A subprocess was required because the failure could be a deadlock

An ordinary numerical regression eventually returns control to Python.

An assertion can then report the mismatch.

This race might never return.

Consider:

Warp 0
→ waits on the current barrier generation
→ completes
→ reaches the tail local barrier

Warp 1
→ waits on another generation
→ never completes
→ never reaches the tail barrier

The kernel call can remain blocked indefinitely.

Running such a test directly in the main pytest process could stop the entire test suite.

PR #11334 therefore executed the delayed kernel inside a separate process.

Parent pytest process
        ↓
Start child process
        ↓
Run the delayed PTX kernel
        ↓
Return normally or exceed the timeout

The test set a 15-second process timeout:

TRITON_TEST_PROCESS_TIMEOUT = 15

If the child’s CUDA context became stuck, the parent could still classify the timeout and keep control of the broader test run.

A Python thread would not provide the same isolation.

Threads can share one process and one CUDA runtime context.

If the context becomes unrecoverable, another thread in the same process may not be able to cleanly terminate or report the result.

For a possible nontermination failure, the process boundary is part of the safety contract.

The design was strong, but it did not remain stable on GB200

The delayed-counter test had several strong properties:

It identified the exact generated instruction boundary.

It separated the first warp from the remaining warps.

It verified that the modified PTX was used.

It isolated a possible deadlock in another process.

It imposed a finite deadline on nontermination.

But a good counterexample design and a stable cross-hardware CI test are not the same thing.

PR #11336 states that the delayed-counter test timed out on GB200.

It then reverted the Python test-file changes introduced by #11334.

The exact scope of the follow-up matters.

Confirmed

The dynamic delayed-counter test timed out on a GB200 job.

The Python delayed-counter test was removed.

The pre-existing ordinary cluster-barrier test was restored.

The compiler race fix remained.

The LLVM regression checks remained.

Not confirmed

The public follow-up does not establish the precise reason for the timeout.

It does not determine whether:

the 10,000,000-clock delay was too long on GB200

the ir_override path behaved differently

the subprocess deadline was too short

the PTX pattern or injected code interacted differently with GB200

another synchronization issue remained

the timeout was caused by a transient CI condition

PR #11336 removed the unstable test.

It did not publish a complete root-cause analysis of the GB200 timeout.

Its validation statement also says that the Python file matched its pre-#11334 form and that repository pre-commit checks passed.

GPU tests were not rerun as part of that follow-up.

The accurate conclusion is therefore:

The PTX-delay method attempted to force the race deterministically, but it did not remain stable in GB200 CI and was removed. The public record examined here does not establish the exact cause of that timeout.

Removing the dynamic test did not revert the compiler fix

The removal of the test does not mean the original race was rejected or that the ordering change was undone.

The current lowering still follows:

load the current counter
→ local barrier
→ derive slot and parity
→ perform the cluster barrier
→ store the next counter
→ tail local barrier

The current source also retains the comment explaining the intended contract:

The warp-specialization entry or the previous tail barrier
publishes the current counter.

Wait for every local thread to read it
before the leader can advance it.

The active code loads the shared counter before creating the local barrier.

After the mbarrier wait, thread 0 stores the next counter and the lowering emits another local barrier.

The two state boundaries are:

Current state consumption:
load → barrier

Next state publication:
store → barrier

The dynamic test was one method of exercising that contract.

The contract itself remains in the compiler.

The static LLVM regression checks still pin the instruction order

The current LLVM lowering test explicitly checks the order around the current counter.

Conceptually:

There must not be a local barrier before the counter load.

The counter load must be generated.

The next instruction must be the local barrier.

At the end of the round, it checks:

the next counter is calculated

the shared counter is stored

the next instruction is the tail local barrier

The retained static contract is therefore:

Current counter:
load → barrier

Next counter:
store → barrier

Equivalent checks remain for the Rubin lowering path as well.

A future compiler change that restores:

barrier → load

should fail this static regression test.

That is meaningful protection.

It is not identical to executing a late-reader schedule on real hardware.

What the retained ordinary GPU test proves

PR #11336 restored the pre-existing Python functional test.

That test runs the cluster-barrier kernel on Hopper-or-later hardware with:

num_ctas = 2

num_ctas = 4

It verifies the generated PTX properties and the final output.

This supports:

The current compiler output executes on the tested GPU path.

The basic warp-specialized cluster barrier works for two- and four-CTA clusters.

The output is correct under the observed execution schedule.

It does not deliberately construct:

Warp 0 reads the counter immediately.

Warps 1–3 remain delayed immediately before the counter load.

The current coverage can therefore be separated as follows:

Ordinary hardware execution
→ dynamic functional test retained

Compiler instruction ordering
→ static LLVM test retained

Forced late-warp scheduling at the counter load
→ dynamic test no longer retained

Static and dynamic tests answer different questions

Static LLVM check

It asks:

Did the compiler generate the required instruction order?

Its strengths are:

fast

deterministic

not dependent on warp scheduling

less sensitive to GPU clock behavior

precise about the reordering being protected

Its limitation is that it does not execute the problematic schedule on hardware.

Ordinary GPU functional test

It asks:

Does the current generated kernel complete and produce the right result on a supported GPU?

Its strengths are:

uses real hardware

passes through the cluster execution path

verifies generated PTX properties

checks final output

Its limitation is that it does not force the specific late-reader counterexample.

Removed PTX delayed-counter test

It asked:

If some local warps are delayed immediately before the counter load, can the first warp advance the shared phase before they have read it?

Its strengths were:

directly targeted the race boundary

attempted to manufacture the dangerous schedule

tested possible nontermination

isolated the failure with a process timeout

Its limitations were:

dependence on generated PTX structure

dependence on clock-based delay behavior

sensitivity to architecture and scheduling

failure to remain stable in GB200 CI

No single test answered every question.

What remains and what was removed

Compiler race fix
→ retained

load → barrier ordering
→ retained

store → barrier tail ordering
→ retained

LLVM static regression checks
→ retained

Ordinary Hopper-or-later GPU functional test
→ retained

PTX clock64 late-warp perturbation
→ removed

Timeout-bounded dynamic delayed-counter test
→ removed

Exact root cause of the GB200 timeout
→ not established in the examined public record

Two distinctions must remain visible:

The dynamic test was removed
≠
the compiler fix was removed

and:

The compiler fix remains
≠
the late-reader schedule is still dynamically tested

The test perturbation had its own hardware contract

A concurrency regression test can itself become low-level systems code.

The removed test depended on several assumptions:

The generated PTX contained the expected instruction pattern.

The first 32 thread IDs represented the first warp.

A 10,000,000-clock spin created a useful but bounded delay.

The fixed kernel completed within the process timeout.

ir_override behaved consistently across the tested targets.

Each assumption can be reasonable.

The combination may still behave differently across GPU generations, compiler versions, or CI environments.

The GB200 timeout demonstrates at least one broader lesson:

A test that manufactures a GPU schedule also has an architecture and runtime contract of its own.

The compiler was not the only low-level program under examination.

The reproducer was another low-level program.

Removing an unstable test should be recorded as a coverage change

Removing a test that can stall CI may be necessary.

A permanently timing-out GPU job can block unrelated development.

But removal should not be described only as cleanup.

Two things happened at once:

CI stability was restored

and:

hardware-level coverage of one late-reader schedule was reduced

The current state is best described as:

The compiler ordering remains protected statically, and an ordinary cluster-barrier GPU test remains. The PTX-based dynamic test that deliberately delayed selected warps at the counter-load boundary is no longer part of the retained suite.

What a future dynamic regression would need

A replacement test should begin by separating the cause of the GB200 timeout.

Did the fixed synchronization protocol actually stall?

Did the PTX spin behave differently?

Was the timeout simply too short?

Was the injected code placed at an unexpected location?

Did the override path behave differently on that target?

Only after that boundary is known can the test be redesigned confidently.

A future test would ideally preserve several properties.

It should still target the real race boundary

The essential schedule remains:

Delayed warp
→ has not loaded the counter

Fast warp
→ has already loaded it

The test must still determine whether the fast warp can advance the state before every local reader has captured the current value.

It should depend less on accidental PTX text structure

A compiler instrumentation hook, a test-only lowering mode, or another structured mechanism might be more stable than matching raw generated PTX with a regular expression.

That is a possible future design direction, not a feature confirmed to exist in the current code.

The delay and timeout should have a target-aware contract

A delay that is too short does not create the schedule.

A delay that is too long can cause even the fixed code to exceed a CI deadline.

The relationship between:

delay length

GPU clock behavior

expected kernel completion time

process timeout

needs to be validated across target classes.

A negative control should be explicit

The strongest evidence would compare:

Old ordering
+
the perturbation
→ fails or times out

with:

Fixed ordering
+
the same perturbation
→ completes

That comparison could use a test-only variant rather than depending on historical source checkout.

The important point is to show that the ordering change, rather than some unrelated condition, controls the result.

Possible nontermination must remain isolated

A deadlock-capable kernel should still run behind:

a separate process

a finite timeout

a parent process capable of reporting the result

That was one of the strongest safety boundaries in the original design.

The removed test still has analytical value

The test no longer exists in current main.

Its design remains worth studying.

It attempted to transform the race from:

an occasional scheduler-dependent delay

into:

a precisely located, deliberately manufactured schedule

The sequence was:

Find the generated counter load
        ↓
Delay selected warps at that exact instruction
        ↓
Allow one warp to run ahead
        ↓
Compare what old and fixed ordering permit
        ↓
Isolate possible deadlock in another process

This is stronger than simply running the kernel many times and hoping for an unlucky schedule.

But the follow-up teaches an additional lesson:

A deterministic counterexample and a stable multi-generation CI test are separate engineering achievements.

The first can be conceptually excellent while the second remains unresolved.

The patch and the test had different life cycles

When PR #11334 first landed, three pieces arrived together:

1. Compiler fix

2. LLVM static regression checks

3. PTX delayed-counter dynamic test

After PR #11336, the retained state became:

1. Compiler fix

2. LLVM static regression checks

3. Restored ordinary GPU functional test

Looking only at the initial merge would miss this difference.

The actual state of a system requires tracking:

what was first introduced

what was removed in a follow-up

what remained after both changes

A merged PR is not always the end of the evidence chain.

The current canonical assessment

Confirmed

The local barrier that previously preceded the shared counter load was moved after the load.
The new ordering prevents the leader from advancing until all participating local threads have completed the current counter read.
Triton main retains the load → barrier ordering.
Triton main retains the store → barrier tail-publication ordering.
Static LLVM regression checks protect both instruction boundaries.
An ordinary Hopper-or-later cluster-barrier GPU functional test remains.
The PTX delayed-counter dynamic test introduced in #11334 was removed by #11336.

Not confirmed

The exact cause of the GB200 timeout is not established by the examined public material.
The timeout does not prove that the retained compiler fix is incorrect.
The timeout also does not prove that every late-reader hardware schedule is currently covered.
The dynamic PTX-delay method is no longer present to provide that specific coverage.

The test for the race became another object of verification

The original target was compiler synchronization.

The PTX test was introduced as the evidence-producing instrument.

The GB200 timeout then made the instrument itself a new boundary to investigate.

Compiler synchronization
→ original subject

PTX perturbation
→ reproduction instrument

GB200 timeout
→ new evidence about the instrument’s limits

A low-level regression test does not automatically produce truth merely because it manipulates PTX.

The test must also be checked for:

whether it modified the intended instruction,
whether the modified code was actually executed,
whether its timing assumptions remain meaningful on the target,
whether a timeout indicates the product defect or the test mechanism.

The same fix now has a different evidence level

The active compiler code is clear:

counter = load();
barrier();

The retained static regression contract is also clear:

The counter load must be immediately followed by the barrier.

The next-counter store must be immediately followed by the tail barrier.

What no longer exists is the original hardware scheduling perturbation.

The final statement should therefore remain precise:

Triton fixed the phase-snapshot ordering in cluster_barrier, and the ordering remains protected by static LLVM regression checks. The PTX delayed-counter test that attempted to exercise the race dynamically was removed after timing out on GB200, so hardware-level regression coverage through that particular late-warp perturbation is no longer retained.

Triton #11334 — Why cluster_barrier Could Advance the Phase Before Every Warp Had Read It
Triton #11334 — How Counter Values 0, 1, 2, and 3 Encode Two mbarrier Slots and Parity

Related material

Compiler fix status: Retained in Triton main
Dynamic PTX regression: Added in PR #11334 and removed by PR #11336 after a GB200 timeout
Current retained coverage: LLVM static instruction-order checks and the ordinary Hopper-or-later GPU functional test
Current evidence gap: No retained dynamic regression currently forces selected warps to remain delayed immediately before the shared counter load

This is Part 3 and the final article in the Triton NVIDIA cluster_barrier series.

Part 1 examined why the old barrier → load ordering did not prove that every warp had completed its phase read.

Part 2 showed how a two-bit counter selected two physical mbarrier slots and parity, and how one late load could split one source-level call across different barrier generations.

This final article examined the generated-PTX delay test that attempted to force the race deterministically, why it was removed after a GB200 timeout, and which regression evidence remains today.

#Triton #NVIDIA #CUDA #Hopper #Blackwell #GB200 #PTX #MBarrier #WarpSpecialization #RaceCondition #TestCoverage #CodeAnalysis