Repository navigation
fix(ck-tile): split the TDM dim1 tensor stride at the right bit - #12871
Merged
Merged
Conversation
TDM_GROUP1::tensorDim1Stride split the value at bit 32 and masked with
0xFFFFFFFF. The field does not split there:
uint32_t tensor_dim1_stride_lo : 16; // high half of sgpr6
uint32_t tensor_dim1_stride_hi; // all of sgpr7
so the split point is bit 16. `lo` survived only by accidental truncation
into its 16-bit bitfield, and `hi` received bits [32:64], silently dropping
bits [16:32] of the stride. The 0xFFFFFFFF mask on a 16-bit field was itself
wrong and helped disguise the shift.
dim1 is the only field in the descriptor where `lo` is narrower than `hi`.
Everywhere else `lo >= hi` and the symmetric mask/shift is correct --
tensorDim0Stride sits directly above this one and is right -- which is what
makes the defect read as correct code.
Also widens tensorDimStride's `value` parameter to uint64_t. Its only caller
passes a uint64_t stride, so the narrow parameter truncated every stride to
32 bits before either setter could place it, making the wider halves of both
stride fields unreachable.
Effect: strides below 2^16 were already correct; at and above 2^16 the
descriptor pointed the DMA engine at an out-of-range address. That surfaces
as a runtime memory fault rather than a wrong result, so no static check
catches it.
Adds test_tdm_descriptor_stride, a descriptor readback that packs strides
through the real setters on device, copies the raw D# words back, and
reconstructs each stride from the bitfield widths. The decoder is written
from the field layout rather than as the setters' inverse, so a matching bug
in both cannot cancel out. dim0 is carried as a control.
Verified on gfx1250: before this change the dim1 case fails for every stride
at or above 2^16 while the dim0 control passes; after it, both pass across
the full 48-bit range. A passing GEMM would not have been evidence here --
the TDM GEMM path has independent problems on this part, so "it works now"
was not reachable.
✅ 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. |
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
The regression test bypasses the widened production dispatcher, leaving part of the fix uncovered.
Review effort: Balanced
Findings: 1
What changed in this PR
Fixes 48-bit TDM stride encoding and adds regression coverage.
Changes:
- Corrects dim1’s split point and preserves 64-bit stride inputs.
- Adds and registers descriptor readback tests.
| File | Description |
|---|---|
amd_tdm_descriptor.hpp |
Corrects stride encoding. |
tdm_descriptor_stride.cpp |
Tests descriptor stride packing. |
CMakeLists.txt |
Registers the new test. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
JiaLuo-CAN
approved these changes
Oct 2, 2026
JiaLuo-CAN
left a comment
JiaLuo-CAN
left a comment
Contributor
There was a problem hiding this comment.
LGTM. Please check the minor change requested by Copilot. Thanks.
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
make_kernel() only uses the functor's type; the call lives in kentry(),
whose body is `Kernel{}(args...)` only under __HIP_DEVICE_COMPILE__. The
host pass therefore never odr-uses PackStridesKernel::operator(), and the
enclosing anonymous namespace gives it internal linkage, which is exactly
when clang emits -Wunused-member-function. The test targets build with
-Weverything -Werror, so that failed the gfx1250 job.
Mark the member [[maybe_unused]], as DecodeLayoutProbe already does in
test/ck_tile/fmha/test_fmha_bwd_tdm_layout.cpp.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Oct 8, 2026
fix(ck-tile): split the TDM dim1 tensor stride at the right
bit (#12871)
## Motivation
This PR fixes a bug in TDM descriptor that affects cases with stride
values > 65K.
## Technical Details
TDM_GROUP1::tensorDim1Stride split the value at bit 32 and masked with
0xFFFFFFFF. The field does not split there:
uint32_t tensor_dim1_stride_lo : 16; // high half of sgpr6
uint32_t tensor_dim1_stride_hi; // all of sgpr7
so the split point is bit 16. `lo` survived only by accidental
truncation into its 16-bit bitfield, and `hi` received bits [32:64],
silently dropping bits [16:32] of the stride. The 0xFFFFFFFF mask on a
16-bit field was itself wrong and helped disguise the shift.
dim1 is the only field in the descriptor where `lo` is narrower than
`hi`. Everywhere else `lo >= hi` and the symmetric mask/shift is correct
-- tensorDim0Stride sits directly above this one and is right -- which
is what makes the defect read as correct code.
Also widens tensorDimStride's `value` parameter to uint64_t. Its only
caller passes a uint64_t stride, so the narrow parameter truncated every
stride to 32 bits before either setter could place it, making the wider
halves of both stride fields unreachable.
Effect: strides below 2^16 were already correct; at and above 2^16 the
descriptor pointed the DMA engine at an out-of-range address. That
surfaces as a runtime memory fault rather than a wrong result, so no
static check catches it.
Adds test_tdm_descriptor_stride, a descriptor readback that packs
strides through the real setters on device, copies the raw D# words
back, and reconstructs each stride from the bitfield widths. The decoder
is written from the field layout rather than as the setters' inverse, so
a matching bug in both cannot cancel out. dim0 is carried as a control.
Verified on gfx1250: before this change the dim1 case fails for every
stride at or above 2^16 while the dim0 control passes; after it, both
pass across the full 48-bit range. A passing GEMM would not have been
evidence here -- the TDM GEMM path has independent problems on this
part, so "it works now" was not reachable.
## Test Plan
Tests in test/ck_tile/tdm/ should pass.
## Test Result
Tests are passing.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
JIRA ID : AICK-2460
jefyang1
added a commit
that referenced
this pull request
Oct 9, 2026
Cross-checked the TDM descriptor against CK's amd_tdm_descriptor.hpp. The field
layout matches exactly -- all 12 TDM_GROUP0 fields and all 20 TDM_GROUP1 fields,
position and width. One value split does not.
tensor_dim1_stride is 48 bits as lo(16) + hi(32): 16 bits in the high half of
sgpr6 plus all of sgpr7. We split it at bit 32, so `hi` received bits [32, 64)
and bits [16, 48) of the stride were dropped. `lo` was right only because a
16-bit field truncates whatever it is given.
intended 0x10000 -> encoded 0x0
intended 0x12345 -> encoded 0x2345
intended 0x100000 -> encoded 0x0
dim1 is the only field in the descriptor where `lo` is narrower than `hi`.
Everywhere else lo >= hi and the symmetric shift is correct -- dim0 sits
directly above it, splits at 32, and is right -- which is what makes the wrong
one read as correct code.
Not reachable before this commit. TDM_MAX_DIM1_STRIDE capped the stride at
0xFFFF, which is exactly the range where both splits agree, and the GEMM wiring
only ever passes 1. Emitted IR is byte-identical to the previous commit for
every TDM parity config; the gate cannot see this change, which is the point of
the new tests.
The cap itself was the deeper error. Both engines carried a comment saying "CK
truncates the dim-1 stride to a 16-bit field", and 827a771's message called
mirroring it deliberate: "Mirrored verbatim rather than corrected -- diverging
from CK's bit layout would be a silent correctness hazard." CK does not
truncate. The field is 48 bits and always was; the old CK code just could not
reach past 16, and CK fixed its own split in #12871 on
2026-10-08. Matching the old behaviour is now the divergence, so the cap is
raised to the real 48-bit bound and dim0 gains the symmetric guard it never had.
Also corrects a units claim while the file is open: the docstring said strides
are "in elements". That holds for dim0. dim1 is in units of dim0_stride, per the
gfx1250 SPG note CK carries in buffer_view.hpp. At rank 2 the only sensible
value is 1, under which both readings agree, so nothing downstream changes --
but it would mislead anyone generalising past rank 2.
Tests, in both engines, mutation-checked by reverting the shift: 4 Python
assertions and 8 C++ checks fail against the old code. They assert the split
point and the lo|hi<<16 round-trip at 0x10000, 0x12345, 2**20 and 2**48-1,
assert dim0's opposite split, and assert that the two halves sit side by side in
word 6 without colliding when both strides are maximal. Every pre-existing test
passed under the defect because none used a stride above 0xFFFF -- the cap saw
to that.
Gate: 69 GREEN / 1 DRIFT at llvm23, the DRIFT being the pre-existing
gfx942_attention_tiled_3d cfg[3]. gemm configs=18, gfx1250_lowering configs=31.
ctest 14/14, host tests 18 in test_tdm (15 before). The two host-suite failures
are byte-identical to HEAD.
This one is verifiable without a GPU -- it is descriptor packing, not emission --
so unlike the TDM wiring it carries no silicon debt.
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.
Motivation
This PR fixes a bug in TDM descriptor that affects cases with stride values > 65K.
Technical Details
TDM_GROUP1::tensorDim1Stride split the value at bit 32 and masked with 0xFFFFFFFF. The field does not split there:
so the split point is bit 16.
losurvived only by accidental truncation into its 16-bit bitfield, andhireceived bits [32:64], silently dropping bits [16:32] of the stride. The 0xFFFFFFFF mask on a 16-bit field was itself wrong and helped disguise the shift.dim1 is the only field in the descriptor where
lois narrower thanhi. Everywhere elselo >= hiand the symmetric mask/shift is correct -- tensorDim0Stride sits directly above this one and is right -- which is what makes the defect read as correct code.Also widens tensorDimStride's
valueparameter to uint64_t. Its only caller passes a uint64_t stride, so the narrow parameter truncated every stride to 32 bits before either setter could place it, making the wider halves of both stride fields unreachable.Effect: strides below 2^16 were already correct; at and above 2^16 the descriptor pointed the DMA engine at an out-of-range address. That surfaces as a runtime memory fault rather than a wrong result, so no static check catches it.
Adds test_tdm_descriptor_stride, a descriptor readback that packs strides through the real setters on device, copies the raw D# words back, and reconstructs each stride from the bitfield widths. The decoder is written from the field layout rather than as the setters' inverse, so a matching bug in both cannot cancel out. dim0 is carried as a control.
Verified on gfx1250: before this change the dim1 case fails for every stride at or above 2^16 while the dim0 control passes; after it, both pass across the full 48-bit range. A passing GEMM would not have been evidence here -- the TDM GEMM path has independent problems on this part, so "it works now" was not reachable.
Test Plan
Tests in test/ck_tile/tdm/ should pass.
Test Result
Tests are passing.
Submission Checklist
JIRA ID : AICK-2460