Add L9/L10/L11 block-distributed transforms (HIP/MFMA/rocWMMA)#7
Open
ahurta92 wants to merge 2 commits into
Open
Add L9/L10/L11 block-distributed transforms (HIP/MFMA/rocWMMA)#7ahurta92 wants to merge 2 commits into
ahurta92 wants to merge 2 commits into
Conversation
Ports the blocked-transform contributions from the blocked-transform branch onto upstream/main without disturbing existing L1-L8 dispatch. New levels share the per-wave corner-turn algorithm but differ in the inner GEMM: L9 (blocked) v_mfma_f64_16x16x4f64 manual fragments, K=16 L10 (blocked-rocwmma) rocWMMA mma_sync, K=16 L11 (blocked-k20) hybrid 16x16x4 + 4x4x4 MFMA, K=20 All three are gated behind __HIP__ in transformbench.cu and validate_levels.hip so the CUDA executable still compiles. The header bodies are likewise wrapped. Validated bit-exact against L1 on MI210 for K=16 (L9, L10) and K=20 (L11). Also adds: validate.hip GPU L1 vs CPU mTxm reference (mirrors transform3d.cc) BLOCKED_TRANSFORM.md design notes, MFMA fragment layouts, perf tables Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Contributor
Author
|
This replaces the previous PR. I'll close the other. |
Closed
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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
1 participant
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.
Ports the blocked-transform contributions from the blocked-transform branch onto upstream/main without disturbing existing L1-L8 dispatch. New levels share the per-wave corner-turn algorithm but differ in the inner GEMM:
L9 (blocked) v_mfma_f64_16x16x4f64 manual fragments, K=16
L10 (blocked-rocwmma) rocWMMA mma_sync, K=16
L11 (blocked-k20) hybrid 16x16x4 + 4x4x4 MFMA, K=20
All three are gated behind HIP in transformbench.cu and validate_levels.hip so the CUDA executable still compiles. The header bodies are likewise wrapped. Validated bit-exact against L1 on MI210 for K=16 (L9, L10) and K=20 (L11).
Also adds:
validate.hip GPU L1 vs CPU mTxm reference (mirrors transform3d.cc)
BLOCKED_TRANSFORM.md design notes, MFMA fragment layouts, perf tables