Triton PR #11765: Why s_barrier Let the Next K Tile Overwrite LDS Reads
Triton PR #11765 fixes a memory-ordering race in the AMD two-cluster BlockPingpong schedule.
The schedule had already issued local-data-share reads for the current matrix tile when it reached a raw s_barrier. The barrier synchronized the participating waves, but the generated ISA could still have outstanding ds_read operations at that point.
Another warp group was then free to write the next K tile into the same LDS buffer. Some of the first group’s reads could observe that newer data instead of the tile they were supposed to consume.
On the reported gfx942 configuration, the corrupted values came from exactly one BLOCK_K ahead:
expected A[r, k]
observed A[r, k + 64]The patch replaces the raw barrier with Triton’s local-memory cluster barrier so the synchronization point carries the required LDS memory-fence semantics.
BlockPingpong overlaps memory work with matrix work
Triton’s AMD BlockPingpong pass reorganizes a GEMM loop so different warp groups alternate between two kinds of instruction clusters:
memory cluster
→ global loads, LDS loads, and LDS stores
dot cluster
→ MFMA matrix instructionsWhile one group performs matrix work, another can prepare data. Priorities and conditional barriers keep the groups moving through the alternating schedule without both entering the same expensive cluster at once.
The word “cluster” here describes a scheduled group of instructions. It is not a multi-CTA hardware cluster.
The two-cluster variant is used for tile shapes that do not fit the pass’s one-cluster or four-cluster schedules. It slices the local loads and dot operation into two parts, then arranges them approximately as:
load first A/B slices from LDS
↓
global load
↓
load second A/B slices from LDS
↓
global load
↓
synchronize
↓
execute first half of the dotThat synchronization point must protect the LDS data just read from being overwritten by the group preparing the next iteration.
The old code used two different kinds of barrier
Most synchronization points in the pass used a helper named appendClusterBarrier().
That helper emits:
ttg.barrier local
rocdl.sched.barrier noneThe first operation carries local-memory synchronization semantics. The second tells the backend not to move scheduled instructions across the boundary.
The first memory cluster in transformTwoPPClusters() was different. It ended with:
appendOp(ROCDL::SBarrierOp::create(builder, loc));
appendOp(ROCDL::SchedBarrier::create(
builder, loc, SchedGroupMask::none));A comment said the cluster was too full to include the wait associated with the local barrier and assumed the backend could place local-memory fences later in the first dot cluster.
That placement was too late.
The memory dependency had to be complete before the other warp group passed the synchronization point and reused the shared buffer.
s_barrier arrived before the LDS reads completed
The issue reporter inspected the generated gfx942 ISA and found this sequence:
ds_read2_b32 ...
buffer_load_dwordx4 ...
s_barrier
s_waitcnt lgkmcnt(14)
v_mfma_f32_32x32x2_f32 ...The ds_read2_b32 instructions read fragments from LDS into registers.
lgkmcnt tracks outstanding local and scalar-memory operations. A partial wait after the barrier did not establish that all LDS reads had completed before another wave moved beyond the same barrier.
The paired barrier later in the loop had a full:
s_waitcnt lgkmcnt(0)before s_barrier.
The failing barrier did not.
A rendezvous placed before the required memory completion cannot protect data that is still being read.
Once the other group passed the barrier, its next-iteration ds_write operations could update the same LDS region. The first group’s unfinished reads then raced with those writes.
The wrong result identified the winning writer
The failure was not random numerical noise.
The issue reporter changed the inputs so every matrix value identified its source position:
B = identity
A[r, k] = 1000 × r + kWith that setup, every corrupted output examined matched:
A[r, k + 64]rather than the expected current-tile value.
The offset was exactly the configured:
BLOCK_K = 64That is strong causal evidence. The affected accumulator fragment consumed the next K-step’s A tile after it had overwritten data still needed from the current step.
The visible pattern was also structured. The report describes 4-row by 32-column patches within the affected 128-by-128 output tiles, concentrated in a specific MFMA accumulator sub-fragment.
A rounding problem would not naturally replace every wrong operand with the corresponding value from precisely one future K tile.
Why another kernel with the same hazard could pass
A tutorial-style matmul using the same main configuration reportedly produced correct results even though its loop had the same missing wait.
Its instruction placement differed.
The failing variants placed the ds_read2 instructions close together immediately before the barrier. The passing variant interleaved pointer-update arithmetic among those reads. That spacing gave more LDS operations time to complete before the other warp group reached its writes.
The missing ordering rule still existed. The timing only determined whether the race became observable.
This is a common property of synchronization defects:
incorrect schedule
+
favorable instruction spacing
=
apparently correct runChanging unrelated arithmetic, compiler scheduling, or register pressure can expose or hide the same underlying race.
A passing timing variant is not proof that the synchronization contract is correct.
The fix moves the memory guarantee to the boundary
The merged patch replaces the special raw barrier with the same helper used at the other cluster boundaries:
// Ensure all local loads complete before the other warp group can overwrite
// the shared-memory buffer for the next iteration.
appendClusterBarrier(builder, loc);This emits a high-level:
ttg.barrier localfollowed by the scheduler barrier.
Triton’s AMD barrier-lowering test shows that a local barrier on gfx942 becomes:
workgroup-scoped release fence for local memory
↓
rocdl.s.barrier
↓
workgroup-scoped acquire fence for local memoryThe fences carry an AMD local-address-space tag so the backend knows which memory domain is being synchronized.
The old raw s_barrier expressed wave rendezvous directly. The new operation states the higher-level requirement: local-memory accesses before the barrier must be ordered against accesses after it.
That gives later lowering enough semantic information to place the required wait at the synchronization boundary rather than after it.
The patch preserves BlockPingpong
The fix does not disable the optimization or route gfx942 away from the two-cluster schedule.
It changes one boundary inside that schedule.
The surrounding structure remains:
interleaved local and global loads
↓
local-memory cluster barrier
↓
first dot slice
↓
cluster barrier
↓
next local stores
↓
cluster barrier
↓
second dot sliceThe pass can still overlap memory and MFMA work. It simply may not exchange ownership of the reusable LDS data until the current local reads are complete.
No performance benchmark was published with the pull request. The added fence can affect scheduling on the targeted path, but the previous schedule produced incorrect arithmetic and could not be preserved as a valid optimization.
The committed test checks compiler structure
The pull request updates five FileCheck expectations in the existing AMD BlockPingpong MLIR test.
Each affected schedule previously expected:
rocdl.s.barrierand now expects:
ttg.barrier localThe surrounding checks still require the sliced local loads, global load, scheduler barrier, priority change, and dot instructions in the intended order.
This protects the compiler-level decision that fixed the bug.
It does not run the reported 256-by-256 matrix multiplication, inspect final s_waitcnt placement, or compare GPU output with torch.mm.
The numerical and ISA evidence comes from issue #11696. The committed regression is structural rather than end to end.
The reported scope is narrow but the rule is general
The failure was observed on gfx942 with:
BLOCK_M = 128
BLOCK_N = 128
BLOCK_K = 64
num_warps = 8
num_stages = 2
kpack = 1
FP32 MFMA 32×32×2The reporter found that only one of the 36 tested PyTorch ROCm matmul configurations for that shape selected the failing two-cluster variant.
Disabling BlockPingpong removed the failure. A gfx950 run did not use this schedule because it took an async-copy path. A different MFMA non-K dimension also avoided the reported failure.
Those observations define where the bug was seen. They do not make the missing memory ordering valid elsewhere.
The general rule is broader:
A synchronization point that permits another wave to reuse shared memory must wait for the current wave’s outstanding shared-memory accesses.
Triton PR #11765 restores that rule at the first boundary of the AMD two-cluster BlockPingpong schedule.
Sources
https://github.com/triton-lang/triton/pull/11765
https://github.com/triton-lang/triton/issues/11696
https://github.com/triton-lang/triton/commit/fe56d08b694285ad4ff923feaa2fdd2792ecb94d