Repository navigation
Conversation
… silently UniversalGemmKernel::IsSupportedArgument rejects any D tensor whose layout differs from CLayout, and is the only rejection in that function that says nothing: every other one emits CK_TILE_ERROR under CK_TILE_LOGGING. The kernel simply never runs, and the reason has to be found by reading the source -- we lost an afternoon to it. The case that hits this is not exotic: a broadcast scale D (an M-vector with a column layout, paired with a row-major C), i.e. per-token scales expressed as a D tensor. Match the surrounding convention and name the offending D index. Compile-checked with hipcc --offload-arch=gfx950 against the unpatched header as a baseline. Co-Authored-By: Claude Fable 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. |
assistant-librarian
Bot
added
the
external contribution
|
Addressed the unit-test policy feedback in
The policy table now reports Unit Test: Pass. The remaining |
|
Hi Thomas, thanks for the pr. Unfortunately, the compilation for gfx11 fails. Please fix. |
|
@illsilin Fixed the gfx11 compilation issue in bb9872f. The synthetic argument-validation fixture models the XDL pipeline interface, so its test target is now registered inside the existing gfx94/gfx95 block instead of the outer block that also builds gfx11/gfx12 WMMA tests. The production header change remains compiled through the real WMMA UniversalGemmKernel users on gfx11/gfx12. |
|
@illsilin Fixed the gfx11 compilation blocker in bb9872f. The new synthetic argument-validation fixture models the XDL pipeline interface, but it was registered outside the existing architecture split and therefore compiled in the gfx11/gfx12 WMMA build. Its target is now scoped to the existing gfx94/gfx95 XDL block. The production header change remains compiled by the real WMMA kernel targets; only this synthetic fixture is excluded from WMMA builds. |
|
Hey Thomas!!!!!! |
|
The new error identifies which D tensor failed, but it does not show the actual D layout or the required C layout. Could we include |
|
Addressed both outstanding review comments in 4df7071:
|
|
@illsilin Could you take another look? |
|
Compilation for gfx950 fails. |
|
@illsilin Hi Illia, Could you share the errors log in the CI? I cannot see our CI dashboard anymore. |
…ow reads UniversalGemmKernel::IsLargeTensorGlobalLoadSupported reads BlockGemmShape::PermuteB and GemmPipeline::Preshuffle, neither of which the mock types declared, so instantiating IsSupportedArgument failed to compile. Add both, matching the defaults in TileGemmShape and GemmPipelineProblem. Mark kclusterM/N/K and BlockSize [[maybe_unused]]: they are declared to mirror the real types but never read, which -Weverything -Werror rejects for an anonymous-namespace type. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
fix(ck_tile): report the D-tensor layout rejection instead of failing silently (#11701) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ISSUE ID : #3766 > Ported from #3768 because that repository is now a read-only mirror. This PR applies the same change to the canonical `projects/composablekernel` tree. Addresses the first ask in #3766. `UniversalGemmKernel::IsSupportedArgument` rejects any D tensor whose layout differs from `CLayout`: ```cpp if(std::is_same_v<DiLayout, CLayout> == false) { DTensorIsValid = false; // no message } ``` It is the **only** rejection in that function that says nothing — every other one emits `CK_TILE_ERROR` under `CK_TILE_LOGGING`. The kernel simply never runs, and the reason has to be found by reading the source. We lost an afternoon to it. The case that hits this is not exotic: a **broadcast scale D** — an M-vector with a column layout, paired with a row-major C — i.e. per-token scales expressed as a D tensor. This PR only makes the rejection speak, matching the surrounding convention and naming the offending D index. The deeper asks in #3766 (support column-layout D in the CShuffle epilogue's distribution, or document RowColQuant as the intended route) are unchanged and left open. Compile-checked with `hipcc -fsyntax-only --offload-arch=gfx950`, with the unpatched header as a baseline. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
TheRock Submodule Bump ActivityNewest first
|
ISSUE ID : ROCm/composable_kernel#3766
Addresses the first ask in ROCm/composable_kernel#3766.
UniversalGemmKernel::IsSupportedArgumentrejects any D tensor whose layout differs fromCLayout:It is the only rejection in that function that says nothing — every other one emits
CK_TILE_ERRORunderCK_TILE_LOGGING. The kernel simply never runs, and the reason has to be found by reading the source. We lost an afternoon to it.The case that hits this is not exotic: a broadcast scale D — an M-vector with a column layout, paired with a row-major C — i.e. per-token scales expressed as a D tensor.
This PR only makes the rejection speak, matching the surrounding convention and naming the offending D index. The deeper asks in ROCm/composable_kernel#3766 (support column-layout D in the CShuffle epilogue's distribution, or document RowColQuant as the intended route) are unchanged and left open.
Compile-checked with
hipcc -fsyntax-only --offload-arch=gfx950, with the unpatched header as a baseline.🤖 Generated with Claude Code