Repository navigation
feat: Compatible K when not a multiple of 8 - #11537
Merged
Merged
Conversation
JP-Fernando
added
project: composablekernel
organization: streamhpc
✅ 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. |
JP-Fernando
commented
Sep 1, 2026
JP-Fernando
changed the title
JP-Fernando
force-pushed
the
users/fjimenez/ck/meta_gfx1250_K_dimension
branch
from
September 4, 2026 11:22
1268e6a to
3c6c293
Compare
chris-tsiaousis-hpc
requested changes
Sep 7, 2026
chris-tsiaousis-hpc
left a comment
chris-tsiaousis-hpc
left a comment
Contributor
There was a problem hiding this comment.
I've added comments/suggestions but my biggest concern is that there are no tests for K%8!=0. Please add them.
Contributor
Author
Added new tests. |
1 task
bidlekm
reviewed
Sep 9, 2026
bidlekm
approved these changes
Sep 17, 2026
chris-tsiaousis-hpc
approved these changes
Sep 17, 2026
chris-tsiaousis-hpc
marked this pull request as ready for review
September 18, 2026 11:20
illsilin
approved these changes
Sep 21, 2026
illsilin
enabled auto-merge (squash)
September 21, 2026 14:25
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Sep 21, 2026
feat: Compatible K when not a multiple of 8 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation When the K dimension is not aligned with the vector size of 8 (which is the case on `gfx1250` for 16-bit datatypes), `IsSupportedArgument` is `false`, so the shape is rejected. This is an unnecessary restriction, since the final partial `K` tile can be zero-padded safely. This PR allows supported GEMM configurations to process `K%8 != 0` correctly. ## Technical Details This PR is currently scoped to `gfx1250` architecture; RCR layouts; and the non-preshuffled ComputeV3 pipeline. Update: We added the vector sizes of A/B/C as parameters to `GemmConfig`, as well as a `FixedVectorSize` switch. Should `K` be a multiple of the operand's vector size, the existing GEMM configuration and loading mechanism are kept unchanged; otherwise, the call is re-dispatched through `GemmConfigFixedVectorSize<>`. ## Test Plan We extracted the provided shapes that satisfy `N%128==0; K%8!=0` and tested them through `tile_example_gemm_universal`. We also built and ran `test_ck_tile_gemm_pipeline_compv3_wmma` to check for regressions. ## Test Result The tests were numerically validated on the Conductor `a07-1` machine, which has `gfx1250` architecture. - `tile_example_gemm_universal`: 13/13 passed - `test_ck_tile_gemm_pipeline_compv3_wmma`: 252 passed, 70 skipped, 0 failed ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Closes #11537
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task
illsilin
added a commit
that referenced
this pull request
Sep 22, 2026
## Motivation When the inner inner matrix dimensin is not aligned with the vector loading size, the shape is rejected by `IsSupportedArgument`. To run these shapes, a switch to scalar loading is necessary in these cases. This PR builds upon: [feat: Compatible K when not a multiple of 8](#11537 (comment)) ## Technical Details There are two GEMM families affected: `universal` and `batched`. `Universal` GEMM: `gemm_utils.hpp` already has a shared `GemmConfigBase` which every config inherits from, with `kPadM/N/K, FixedVectorSize, VectorSizeA/B/C` all defaulting to `false/1`. A new config got added which inherits from whichever `GemmConfig` is passed as a template parameter and overrides just the mentioned fields. `run_gemm_example_with_layouts` contains a decision tree: it computes what configuration to pick based on the vector load size and the passed shape. `Batched` GEMM: `batched_gemm.hpp` configs are standalone structs with no shared base class, they never had `kPadM/N/K/FixedVectorSize/VectorSizeA/B/C` fields at all. A new config is added that has these values, but to let `batched_gemm()` run with either the pre-existing configs (which lack them) or the new one, an SFINAE wrapper with field detection (GemmConfigVectorSizeDefaults) was added. This resolves to the config itself if the field is present, or to an all-`false/1` defaults struct otherwise, with the same overriding mechanism as in `Universal` GEMM. Closes: #11776 ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Fernando Jiménez <fernando.jimenez@streamhpc.com> Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Sep 22, 2026
feat: support dimension sizes less than 8 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation When the inner inner matrix dimensin is not aligned with the vector loading size, the shape is rejected by `IsSupportedArgument`. To run these shapes, a switch to scalar loading is necessary in these cases. This PR builds upon: [feat: Compatible K when not a multiple of 8](ROCm/rocm-libraries#11537 (comment)) ## Technical Details There are two GEMM families affected: `universal` and `batched`. `Universal` GEMM: `gemm_utils.hpp` already has a shared `GemmConfigBase` which every config inherits from, with `kPadM/N/K, FixedVectorSize, VectorSizeA/B/C` all defaulting to `false/1`. A new config got added which inherits from whichever `GemmConfig` is passed as a template parameter and overrides just the mentioned fields. `run_gemm_example_with_layouts` contains a decision tree: it computes what configuration to pick based on the vector load size and the passed shape. `Batched` GEMM: `batched_gemm.hpp` configs are standalone structs with no shared base class, they never had `kPadM/N/K/FixedVectorSize/VectorSizeA/B/C` fields at all. A new config is added that has these values, but to let `batched_gemm()` run with either the pre-existing configs (which lack them) or the new one, an SFINAE wrapper with field detection (GemmConfigVectorSizeDefaults) was added. This resolves to the config itself if the field is present, or to an all-`false/1` defaults struct otherwise, with the same overriding mechanism as in `Universal` GEMM. Closes: #11776 ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
illsilin
pushed a commit
that referenced
this pull request
Sep 22, 2026
## Motivation Fix the pipeline failure of gfx90a caused by the use of C++20 feature std::type_identity in PR #11537. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
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
When the K dimension is not aligned with the vector size of 8 (which is the case on
gfx1250for 16-bit datatypes),IsSupportedArgumentisfalse, so the shape is rejected.This is an unnecessary restriction, since the final partial
Ktile can be zero-padded safely. This PR allows supported GEMM configurations to processK%8 != 0correctly.Technical Details
This PR is currently scoped to
gfx1250architecture; RCR layouts; and the non-preshuffled ComputeV3 pipeline.Update: We added the vector sizes of A/B/C as parameters to
GemmConfig, as well as aFixedVectorSizeswitch.Should
Kbe a multiple of the operand's vector size, the existing GEMM configuration and loading mechanism are kept unchanged; otherwise, the call is re-dispatched throughGemmConfigFixedVectorSize<>.Test Plan
We extracted the provided shapes that satisfy
N%128==0; K%8!=0and tested them throughtile_example_gemm_universal. We also built and rantest_ck_tile_gemm_pipeline_compv3_wmmato check for regressions.Test Result
The tests were numerically validated on the Conductor
a07-1machine, which hasgfx1250architecture.tile_example_gemm_universal: 13/13 passedtest_ck_tile_gemm_pipeline_compv3_wmma: 252 passed, 70 skipped, 0 failedSubmission Checklist
Closes #11537