fix(ck): fix test_gemm_mx correctness failures on gfx1250 (LDS read drain) - #10653
Merged
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
andriy-ca
changed the title
andriy-ca
marked this pull request as ready for review
August 12, 2026 03:59
Contributor
|
LGTM. |
Contributor
There was a problem hiding this comment.
Pull request overview
Adds explicit gfx1250 LDS-read draining to prevent asynchronous copies from overwriting buffers still being read.
Changes:
- Waits for
DScntbefore the async-load barrier. - Documents the cross-wave ordering requirement.
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
andriy-ca
enabled auto-merge (squash)
August 12, 2026 18:22
JiaLuo-CAN
approved these changes
Aug 12, 2026
geyyer
approved these changes
Aug 12, 2026
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Aug 13, 2026
fix(ck): fix test_gemm_mx correctness failures on gfx1250
(LDS read drain) (#10653)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
JIRA ID : AICK-1890
## Summary
`test_gemm_mx` fails the correctness check on gfx1250 (A0). Five tests
produce sparse wrong values, non-deterministically, at roughly a 50% hit
rate per run:
```
TestGemmMX_MK_NK/0.Large f8 x f8 -> f16
TestGemmMX_MK_NK/1.Large f8 x f8 -> bf16
TestGemmMX_MK_NK/2.Large f4 x f4 -> f16
TestGemmMX_MK_NK/3.Large f6 x f6 -> f16
TestGemmMX_MK_NK/4.Large bf6 x bf6 -> bf16
```
The MX v3 pipeline double-buffers LDS. Each iteration reads one buffer
with `ds_load` while the hardware async copy
`global_load_async_to_lds_b128` fills the other. The barrier at the top
of `LoopFunc` is what keeps a wave from overwriting a buffer another
wave is still reading.
On gfx1250 that barrier is `block_sync_lds_async_load()`, which waits on
**ASYNCcnt**. LDS reads are tracked by **DScnt**, a different counter,
so the barrier does not wait for them. The compiler normally emits the
LDS wait itself, but in this hot loop it computes the weakest wait its
per-wave dependency analysis requires and sinks it *past* the barrier.
In the generated ISA a barrier retires with **28 `ds_load`s still
outstanding** (`s_wait_loadcnt_dscnt 0x11c`), immediately followed by an
async write into that same buffer.
Per-wave dependency analysis cannot see the cross-wave contract — that
the barrier exists so *other* waves may overwrite the buffer — so the
ordering has to be explicit in the source.
## The change
One line, in the existing gfx1250 arm of this one pipeline:
```cpp
#if defined(__gfx125__)
llvm_amdgcn_s_wait_dscnt(0);
block_sync_lds_async_load();
#else
```
A full drain the compiler cannot move past the barrier.
## Test plan
All runs on gfx1250 (ASIC rev 0x0).
- [x] `test_gemm_mx` full suite: **5 failing tests on every run → 0
failures, 3/3 runs**
- [x] Isolated instance, 20 repetitions per build, measured back to
back: **11/20 failures → 0/20**
- [x] `example_gemm_mx_fp8` (configured identically to the failing
instance): **10/10 incorrect → 0/10**
- [x] ISA verified: the added drain appears before the barrier at every
async site; `ds_load` count unchanged, so this is an ordering fix and
not a data-flow change
- [x] Non-gfx1250 paths untouched (change is inside `#if
defined(__gfx125__)`)
Reproducer, for anyone verifying:
```
./bin/example_gemm_mx_fp8 1 2 1 0 5120 5120 4096 4096 4096 5120 1 20 50
```
This fails on **every** run before the fix, which makes it a much better
regression gate than the test suite's ~50% hit rate.
## Performance
~1.6% throughput on the affected kernel (351.5 → 345.9 TFlops at
5120x5120x4096). The baseline is computing wrong answers, so this is the
cost of correctness rather than a regression against a working build.
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Aug 18, 2026
fix(ck): fix test_gemm_mx correctness failures on gfx1250 (LDS read drain)
JIRA ID : AICK-1890
## Summary
`test_gemm_mx` fails the correctness check on gfx1250 (A0). Five tests
produce sparse wrong values, non-deterministically, at roughly a 50% hit
rate per run:
```
TestGemmMX_MK_NK/0.Large f8 x f8 -> f16
TestGemmMX_MK_NK/1.Large f8 x f8 -> bf16
TestGemmMX_MK_NK/2.Large f4 x f4 -> f16
TestGemmMX_MK_NK/3.Large f6 x f6 -> f16
TestGemmMX_MK_NK/4.Large bf6 x bf6 -> bf16
```
The MX v3 pipeline double-buffers LDS. Each iteration reads one buffer
with `ds_load` while the hardware async copy
`global_load_async_to_lds_b128` fills the other. The barrier at the top
of `LoopFunc` is what keeps a wave from overwriting a buffer another
wave is still reading.
On gfx1250 that barrier is `block_sync_lds_async_load()`, which waits on
**ASYNCcnt**. LDS reads are tracked by **DScnt**, a different counter,
so the barrier does not wait for them. The compiler normally emits the
LDS wait itself, but in this hot loop it computes the weakest wait its
per-wave dependency analysis requires and sinks it *past* the barrier.
In the generated ISA a barrier retires with **28 `ds_load`s still
outstanding** (`s_wait_loadcnt_dscnt 0x11c`), immediately followed by an
async write into that same buffer.
Per-wave dependency analysis cannot see the cross-wave contract — that
the barrier exists so *other* waves may overwrite the buffer — so the
ordering has to be explicit in the source.
## The change
One line, in the existing gfx1250 arm of this one pipeline:
```cpp
#if defined(__gfx125__)
llvm_amdgcn_s_wait_dscnt(0);
block_sync_lds_async_load();
#else
```
A full drain the compiler cannot move past the barrier.
## Test plan
All runs on gfx1250 (ASIC rev 0x0).
- [x] `test_gemm_mx` full suite: **5 failing tests on every run → 0
failures, 3/3 runs**
- [x] Isolated instance, 20 repetitions per build, measured back to
back: **11/20 failures → 0/20**
- [x] `example_gemm_mx_fp8` (configured identically to the failing
instance): **10/10 incorrect → 0/10**
- [x] ISA verified: the added drain appears before the barrier at every
async site; `ds_load` count unchanged, so this is an ordering fix and
not a data-flow change
- [x] Non-gfx1250 paths untouched (change is inside `#if
defined(__gfx125__)`)
Reproducer, for anyone verifying:
```
./bin/example_gemm_mx_fp8 1 2 1 0 5120 5120 4096 4096 4096 5120 1 20 50
```
This fails on **every** run before the fix, which makes it a much better
regression gate than the test suite's ~50% hit rate.
## Performance
~1.6% throughput on the affected kernel (351.5 → 345.9 TFlops at
5120x5120x4096). The baseline is computing wrong answers, so this is the
cost of correctness rather than a regression against a working build.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AICK-1890
Summary
test_gemm_mxfails the correctness check on gfx1250 (A0). Five tests produce sparse wrong values, non-deterministically, at roughly a 50% hit rate per run:The MX v3 pipeline double-buffers LDS. Each iteration reads one buffer with
ds_loadwhile the hardware async copyglobal_load_async_to_lds_b128fills the other. The barrier at the top ofLoopFuncis what keeps a wave from overwriting a buffer another wave is still reading.On gfx1250 that barrier is
block_sync_lds_async_load(), which waits on ASYNCcnt. LDS reads are tracked by DScnt, a different counter, so the barrier does not wait for them. The compiler normally emits the LDS wait itself, but in this hot loop it computes the weakest wait its per-wave dependency analysis requires and sinks it past the barrier. In the generated ISA a barrier retires with 28ds_loads still outstanding (s_wait_loadcnt_dscnt 0x11c), immediately followed by an async write into that same buffer.Per-wave dependency analysis cannot see the cross-wave contract — that the barrier exists so other waves may overwrite the buffer — so the ordering has to be explicit in the source.
The change
One line, in the existing gfx1250 arm of this one pipeline:
A full drain the compiler cannot move past the barrier.
Test plan
All runs on gfx1250 (ASIC rev 0x0).
test_gemm_mxfull suite: 5 failing tests on every run → 0 failures, 3/3 runsexample_gemm_mx_fp8(configured identically to the failing instance): 10/10 incorrect → 0/10ds_loadcount unchanged, so this is an ordering fix and not a data-flow change#if defined(__gfx125__))Reproducer, for anyone verifying:
This fails on every run before the fix, which makes it a much better regression gate than the test suite's ~50% hit rate.
Performance
~1.6% throughput on the affected kernel (351.5 → 345.9 TFlops at 5120x5120x4096). The baseline is computing wrong answers, so this is the cost of correctness rather than a regression against a working build.