Repository navigation
feat(ck): convolution LDS Transpose and TransposeC support - #11905
Merged
Merged
Conversation
github-actions
Bot
added
project: composablekernel
ck: convolution
EnricoDeg
changed the title
✅ 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. |
chris-tsiaousis-hpc
approved these changes
Sep 17, 2026
chris-tsiaousis-hpc
left a comment
chris-tsiaousis-hpc
left a comment
Contributor
There was a problem hiding this comment.
LGTM! Added some considerations/suggestions.
EnricoDeg
force-pushed
the
users/enricodeg/ck/lds-transpose-wmma
branch
from
September 21, 2026 08:30
a2be2d8 to
f31f966
Compare
bidlekm
approved these changes
Sep 21, 2026
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Sep 23, 2026
feat(ck): convolution LDS Transpose and TransposeC support (#11905) ## Motivation gfx1250 has LDS transpose instructions, so we can avoid registers permutations for convolution bwd data and bwd weight and directly load the transposed data from LDS into registers. Additionally, TransposeC was not supported by WmmaCShuffleV3 implementation. By using TransposeC we can avoid to have in the epilogue multiple `ds_store_b16` instructions and have instead a single `ds_store_b128` instruction for each wave tile. JIRA ID : ROCM-27829 ## Technical Details - Add transpose support in `ThreadwiseTensorSliceTransfer_v4` used for loading from LDS into registers - Add support in `ABTransferThreadTiles` to load data from gmem in MN Major layout and store it still in MN Major layout (no in register transpose between loading and storing) - Create simple MN Major LDS layout with padding based on K Major layout - Create descriptors in the epilogue for TransposeC case - Expose new features in convolution bwd data and bwd weight with `UseLdsTranspose` and `TransposeC` bool template parameters - Add instances for gfx1250 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.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
gfx1250 has LDS transpose instructions, so we can avoid registers permutations for convolution bwd data and bwd weight and directly load the transposed data from LDS into registers. Additionally, TransposeC was not supported by WmmaCShuffleV3 implementation. By using TransposeC we can avoid to have in the epilogue multiple
ds_store_b16instructions and have instead a singleds_store_b128instruction for each wave tile.JIRA ID : ROCM-27829
Technical Details
ThreadwiseTensorSliceTransfer_v4used for loading from LDS into registersABTransferThreadTilesto load data from gmem in MN Major layout and store it still in MN Major layout (no in register transpose between loading and storing)UseLdsTransposeandTransposeCbool template parametersSubmission Checklist