Accept Hopper matmuls and update default heuristic #3579
Conversation
This file contains 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
jacobhinkle
added a commit
that referenced
this pull request
Dec 16, 2024
This enables Hopper matmul in our automatic scheduler by translating them without introducing new broadcasts. Specifically: 1. Update `mma_utils::MatmulPattern::translateToMmaOp` to optionally avoid intermediates by using an `MmaOp::AxisMapping`. Enable this option when the target arch is not Ampere or Turing. 3. Unguard some tests in `test_translate_mma.cpp` This does not update the default heuristic or change the `canSchedule` checks. See #3579 for that follow-up PR --------- Co-authored-by: Ryan Spring <[email protected]> Co-authored-by: Naoya Maruyama <[email protected]> Co-authored-by: Jingyue Wu <[email protected]> Co-authored-by: nsarka <[email protected]> Co-authored-by: Protonu <[email protected]> Co-authored-by: samnordmann <[email protected]>
Must have been a broken merge
I'm still skipping the ones with batch dimensions on A since these hit an error currently. Will investigate later but we only need 2d A for now.
jacobhinkle
commented
Dec 19, 2024
Comment on lines
-1868
to
-1877
| axis_mapping.a_axes.push_back(d); | ||
| } | ||
| axis_mapping.a_axes.reserve(out_dim); | ||
| for (size_t d : c10::irange(out_dim - 2)) { |
There was a problem hiding this comment.
I think this was just due to a busted merge.
jacobhinkle
commented
Dec 19, 2024
| macro_encode.n = 256; | ||
| while (macro_encode.n >= 8) { | ||
| if (n_extent % macro_encode.n != 0) { | ||
| macro_encode.n /= 2; |
There was a problem hiding this comment.
Currently this only chooses powers of two. For small problems I think we could choose one of the other sizes. For example if n_extent == 72 then we should probably use that size.
jacobhinkle
commented
Dec 19, 2024
|
|
||
| const auto tryIncreaseM = [&]() { | ||
| if (ratiosValid(m_ratio + 1, n_ratio)) { | ||
| m_ratio++; |
There was a problem hiding this comment.
Should these also be powers of two? Currently this will chooses sizes like 192
Should fix this for both matmul and linear, and for avoid_intermediates_ and not
jacobhinkle
changed the title
The dtype for stmatrix should have never been constrained to only Half. The only constraint we have is that the dtype size is 16-bit. This PR is needed for us to actually use stmatrix in bfloat16 matmuls.
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.
No description provided.