Repository navigation
AICK-2528: resolve LDS access offsets against the real descriptor - #13370
Merged
Merged
Conversation
test_load_and_convert_tile_no_transpose/6 (fp8 -> fp8) fails on gfx1201. Two independent bugs interact. The LDS fast path in tile_window's load_access_range() and load_with_offset() was added as a gfx125 optimization, but it is reached through an `if constexpr` that only tests the address space and the absence of an xor swizzle, so gfx1200/gfx1201 take it too. It derives lds_access_offset at compile time from a *default constructed* BottomTensorView. The ys index step is genuinely static, but the descriptor's lengths and strides are not: this test builds its LDS view from runtime extents, so the default-constructed descriptor carries zero strides and every access past the first reads the wrong address. Take the offset from the real descriptor instead. The LDS buffer_view sums static_offset into the same address expression as the runtime linear offset, so where the descriptor is static -- every case the optimization was written for -- this is the same computation and codegen is unchanged. The test then masked that. BlockSize is derived from get_warp_size(), which reports wave64 on the host pass, so the host launched 1024 threads while the wave32 kernel only addressed 512. The surplus threads redundantly rewrote the same output with correct values, hiding the corruption except where the two writers raced. Halve the launch on wave32 devices. Both fixes have to land together: the launch-size fix alone removes the masking without correcting the offsets. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 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. |
illsilin
approved these changes
Oct 9, 2026
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Oct 9, 2026
AICK-2528: resolve LDS access offsets against the real
descriptor (#13370)
JIRA ID : AICK-2528
## Motivation
test_load_and_convert_tile_no_transpose/6 (fp8 -> fp8) fails on gfx1201
in TheRock build. CK did not change: a compiler update exposed a latent
defect that had been present and silently masked. Two independent bugs
interact, one corrupting LDS addresses and the other hiding the
corruption.
## Technical Details
### Bug 1: the LDS fast path resolves offsets against a fabricated
descriptor
The LDS fast path in `tile_window`'s `load_access_range()` and
`load_with_offset()` was added as a gfx125 optimization, but it is
reached through an `if constexpr` that only tests the address space and
the absence of an xor swizzle. It has no architecture component, so
gfx1200/gfx1201, and gfx942/gfx950 via the eight-waves pipeline, take it
too.
It derived `lds_access_offset` at compile time from a *default
constructed* `BottomTensorView`:
```cpp
constexpr auto coord_ys_offset = make_tensor_coordinate(
typename Base::BottomTensorView{}.get_tensor_descriptor(), ...);
```
The ys index step is genuinely static, but the descriptor's lengths and
strides are not. `embed` stores `coefficients_` as a real data member,
so default construction zeroes every stride and `get_offset()` returns 0
for *every* access. Accesses 1, 2, 3 and beyond all read access 0's
address.
This is not a case of C++ computing the wrong thing. `typename
Base::BottomTensorView{}` is a perfectly valid constant expression, just
full of zeros. `constexpr` guarantees *a* constant, never a *correct*
one. The fix takes the offset from the real descriptor instead, which
forces `constexpr` to `const index_t` and folds the offset into
`get_vectorized_elements`' runtime `linear_offset` argument rather than
passing it as a non-type template parameter.
The LDS `buffer_view` sums `static_offset` into the same address
expression as the runtime linear offset, and there is no separate
immediate-field mechanism, so where the descriptor is static this is the
same address computation. Production LDS block descriptors are
`constexpr` with `number<>` strides, so that is the path they all take,
and no change in generated code is expected for them.
Only the final ys-index to byte-offset descriptor lookup changes.
`idx_off_ys` and `adapter_ys_offset` remain `constexpr`, so all the
tile-distribution math stays unconditionally compile-time.
### Bug 2: the test masked it
`BlockSize` derives from `get_warp_size()`, which reports wave64 on the
host pass. The host therefore launched 1024 threads while the wave32
kernel only addressed 512. The surplus threads redundantly rewrote the
same output with correct values, so the corruption from Bug 1 was
overwritten before it could be observed.
Which writer lands last is decided by wave scheduling, so the masking
was never guaranteed. It held in practice until a compiler update
retimed the kernel so that the thread writing the correct value now
loses. Halve the launch on wave32 devices.
Both fixes have to land together: the launch-size fix alone removes the
masking without correcting the offsets.
### Alternatives considered
**Guard the fast path to `#if defined(__gfx125__)`.** This encodes
architecture as a proxy for the real precondition (a static descriptor),
leaves the defect armed on gfx1250 where it would eventually bite, and
would *remove* the optimization from gfx950's eight-waves pipeline,
which reaches it unguarded today. That is a larger behavioral change for
CDNA than this patch makes.
**Gate on `is_static()` / `is_known_at_compile_time()`.** This looks
like the exact precondition and is unsound. The trait falls back to
`std::is_arithmetic` for types lacking the member, so a plain runtime
`index_t` reports `true`. Measured directly: `is_static()` returns
`true` for the very descriptor that fails, so the guard would evaluate
true on the failing test and leave the bug in place. No strict
alternative trait exists, and building one would mean threading it
through every transform in `coordinate_transform.hpp`, a much larger
diff touching a trait every descriptor depends on.
Rewriting the offset is correct everywhere and needs no predicate.
## Test Plan
Local verification on gfx1201 isolates the compiler as the exposing
variable: run `test_load_and_convert_tile_no_transpose` unfixed using
the default compiler on gfx1201, unfixed using TheRock's, and fixed
using TheRock's.
Remaining coverage relies on CI:
- **gfx950**, highest priority. The eight-waves pipeline calls
`load_with_offset` on an LDS window with no arch guard, and gfx950 is
the only target that builds `test_load_and_convert_tile_transposed`. Not
tested locally.
- **gfx1250**, the architecture the fast path was written for.
- **gfx942** and **gfx1200**, same generic path.
- **gfx1100/1101**, the other wave32 family, so the `BlockSize()`
halving applies there as well.
- GEMM, gemm_block_scale, flatmm and FMHA suites, plus perf on CDNA.
## Test Result
gfx1201:
| build | CK | runs | result |
|-------|----|------|--------|
| ROCm 7.2 (default) | unfixed | 100 | all pass |
| TheRock | unfixed | 1 | fail |
| TheRock | fixed | 100 | all pass |
The defect is not visible using the default compiler on the machine and
reproduces immediately using TheRock's. With the fix, 100 consecutive
runs pass using the compiler that exposes it.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
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-2528
Motivation
test_load_and_convert_tile_no_transpose/6 (fp8 -> fp8) fails on gfx1201 in TheRock build. CK did not change: a compiler update exposed a latent defect that had been present and silently masked. Two independent bugs interact, one corrupting LDS addresses and the other hiding the corruption.
Technical Details
Bug 1: the LDS fast path resolves offsets against a fabricated descriptor
The LDS fast path in
tile_window'sload_access_range()andload_with_offset()was added as a gfx125 optimization, but it is reached through anif constexprthat only tests the address space and the absence of an xor swizzle. It has no architecture component, so gfx1200/gfx1201, and gfx942/gfx950 via the eight-waves pipeline, take it too.It derived
lds_access_offsetat compile time from a default constructedBottomTensorView:The ys index step is genuinely static, but the descriptor's lengths and strides are not.
embedstorescoefficients_as a real data member, so default construction zeroes every stride andget_offset()returns 0 for every access. Accesses 1, 2, 3 and beyond all read access 0's address.This is not a case of C++ computing the wrong thing.
typename Base::BottomTensorView{}is a perfectly valid constant expression, just full of zeros.constexprguarantees a constant, never a correct one. The fix takes the offset from the real descriptor instead, which forcesconstexprtoconst index_tand folds the offset intoget_vectorized_elements' runtimelinear_offsetargument rather than passing it as a non-type template parameter.The LDS
buffer_viewsumsstatic_offsetinto the same address expression as the runtime linear offset, and there is no separate immediate-field mechanism, so where the descriptor is static this is the same address computation. Production LDS block descriptors areconstexprwithnumber<>strides, so that is the path they all take, and no change in generated code is expected for them.Only the final ys-index to byte-offset descriptor lookup changes.
idx_off_ysandadapter_ys_offsetremainconstexpr, so all the tile-distribution math stays unconditionally compile-time.Bug 2: the test masked it
BlockSizederives fromget_warp_size(), which reports wave64 on the host pass. The host therefore launched 1024 threads while the wave32 kernel only addressed 512. The surplus threads redundantly rewrote the same output with correct values, so the corruption from Bug 1 was overwritten before it could be observed.Which writer lands last is decided by wave scheduling, so the masking was never guaranteed. It held in practice until a compiler update retimed the kernel so that the thread writing the correct value now loses. Halve the launch on wave32 devices.
Both fixes have to land together: the launch-size fix alone removes the masking without correcting the offsets.
Alternatives considered
Guard the fast path to
#if defined(__gfx125__). This encodes architecture as a proxy for the real precondition (a static descriptor), leaves the defect armed on gfx1250 where it would eventually bite, and would remove the optimization from gfx950's eight-waves pipeline, which reaches it unguarded today. That is a larger behavioral change for CDNA than this patch makes.Gate on
is_static()/is_known_at_compile_time(). This looks like the exact precondition and is unsound. The trait falls back tostd::is_arithmeticfor types lacking the member, so a plain runtimeindex_treportstrue. Measured directly:is_static()returnstruefor the very descriptor that fails, so the guard would evaluate true on the failing test and leave the bug in place. No strict alternative trait exists, and building one would mean threading it through every transform incoordinate_transform.hpp, a much larger diff touching a trait every descriptor depends on.Rewriting the offset is correct everywhere and needs no predicate.
Test Plan
Local verification on gfx1201 isolates the compiler as the exposing variable: run
test_load_and_convert_tile_no_transposeunfixed using the default compiler on gfx1201, unfixed using TheRock's, and fixed using TheRock's.Remaining coverage relies on CI:
load_with_offseton an LDS window with no arch guard, and gfx950 is the only target that buildstest_load_and_convert_tile_transposed. Not tested locally.BlockSize()halving applies there as well.Test Result
gfx1201:
The defect is not visible using the default compiler on the machine and reproduces immediately using TheRock's. With the fix, 100 consecutive runs pass using the compiler that exposes it.
Submission Checklist