Repository navigation
feat: support dimension sizes less than 8 - #11776
Merged
Merged
Conversation
bidlekm
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. |
bidlekm
marked this pull request as draft
September 8, 2026 11:42
chris-tsiaousis-hpc
requested review from
JP-Fernando and
chris-tsiaousis-hpc
September 8, 2026 11:46
…s tile-size misaligned but vector-aligned Adds needsPadM/N/K checks (independent of the existing vector-size needsScalarAB/needsScalarC checks) to both the universal and batched GEMM example dispatch, and a matching GemmConfigFixedVectorSize fallback branch using natural (non-scalar) vector sizes with kPadM/N/K forced on. Also adds test_gemm_unaligned_mnk.cpp covering unaligned M/N (in addition to the existing unaligned-K coverage).
JP-Fernando
requested changes
Sep 10, 2026
JP-Fernando
left a comment
JP-Fernando
left a comment
Contributor
There was a problem hiding this comment.
I have got questions on the compatibility of some pipelines, and whether we should ship this work as-is or whether we should look into it a bit further before merging.
chris-tsiaousis-hpc
requested changes
Sep 10, 2026
chris-tsiaousis-hpc
left a comment
chris-tsiaousis-hpc
left a comment
Contributor
There was a problem hiding this comment.
Left some comments and a more unified design proposal...
chris-tsiaousis-hpc
approved these changes
Sep 14, 2026
bidlekm
marked this pull request as ready for review
September 17, 2026 08:29
JP-Fernando
approved these changes
Sep 17, 2026
illsilin
approved these changes
Sep 18, 2026
illsilin
approved these changes
Sep 18, 2026
illsilin
enabled auto-merge (squash)
September 18, 2026 20:00
…OCm/rocm-libraries into users/bidlekm/ck/lower-dim-gemm-enable
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.
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
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
Technical Details
There are two GEMM families affected:
universalandbatched.UniversalGEMM:gemm_utils.hppalready has a sharedGemmConfigBasewhich every config inherits from, withkPadM/N/K, FixedVectorSize, VectorSizeA/B/Call defaulting tofalse/1. A new config got added which inherits from whicheverGemmConfigis passed as a template parameter and overrides just the mentioned fields.run_gemm_example_with_layoutscontains a decision tree: it computes what configuration to pick based on the vector load size and the passed shape.BatchedGEMM:batched_gemm.hppconfigs are standalone structs with no shared base class, they never hadkPadM/N/K/FixedVectorSize/VectorSizeA/B/Cfields at all. A new config is added that has these values, but to letbatched_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/1defaults struct otherwise, with the same overriding mechanism as inUniversalGEMM.Closes: #11776
Test Plan
Test Result
Submission Checklist