Remove LdMatrixTranspose type from LoadStoreOp types#2315
Merged
Conversation
zasdfgbnm
reviewed
May 29, 2024
zasdfgbnm
reviewed
May 29, 2024
Collaborator
Author
|
!build |
naoyam
reviewed
May 29, 2024
zasdfgbnm
reviewed
May 29, 2024
jacobhinkle
reviewed
May 29, 2024
Collaborator
Author
|
!build |
protonu
force-pushed
the
pbasu_remove_ldmatrixtranspose
branch
from
1ce518a to
ddc169f
Compare
Collaborator
Author
|
Terribly sorry for the force push - I had gotten myself into a mess. |
Collaborator
Author
|
!build |
jacobhinkle
reviewed
May 31, 2024
Collaborator
jacobhinkle
left a comment
jacobhinkle
left a comment
There was a problem hiding this comment.
LGTM but I'll leave it to @zasdfgbnm to give the final approval.
zasdfgbnm
approved these changes
May 31, 2024
Collaborator
zasdfgbnm
left a comment
zasdfgbnm
left a comment
There was a problem hiding this comment.
A few comments on the clarity of comment. This code should be working.
Collaborator
Author
|
!build |
zasdfgbnm
approved these changes
May 31, 2024
Co-authored-by: Gao, Xiang <qasdfgtyuiop@gmail.com>
Collaborator
Author
|
!build |
jacobhinkle
reviewed
May 31, 2024
Co-authored-by: Jacob Hinkle <1454944+jacobhinkle@users.noreply.github.com>
Collaborator
Author
|
!build |
protonu
added a commit
that referenced
this pull request
Jun 4, 2024
…eduler (#2309) In this PR we extend the matmul scheduler to support inputs with allocation domains. To the fusion (with inputs tv_a and tv_b), we add two LoadStoreOps to both inputs. The first Op corresponds to a load to shared memory, where we propagate the allocation domain. The second op corresponds to reading to registers, where we don't propagate the allocation domain since the scheduler takes charge of setting the allocation domain in the registers. Based on the difference in the (maybe)allocation domain of the producer and consumer of the second LoadStoreOp, we may do transposed load when reading to registers.  See also #2315.
zasdfgbnm
added a commit
that referenced
this pull request
Jun 5, 2024
This PR modifies how we determine when we do a LdMatrix vs LdMatrixTranspose. We look at the (maybe) allocation domain of the producer and consumer of the load store op do determine if a transpose is required. We get the allocation domain of the consumer and check if the innermost dimension maps to the innermost dimension of the allocation domain of the producer. We transpose if it doesn't map. --------- Co-authored-by: Gao, Xiang <qasdfgtyuiop@gmail.com> Co-authored-by: Jacob Hinkle <1454944+jacobhinkle@users.noreply.github.com>
zasdfgbnm
pushed a commit
that referenced
this pull request
Jun 5, 2024
…eduler (#2309) In this PR we extend the matmul scheduler to support inputs with allocation domains. To the fusion (with inputs tv_a and tv_b), we add two LoadStoreOps to both inputs. The first Op corresponds to a load to shared memory, where we propagate the allocation domain. The second op corresponds to reading to registers, where we don't propagate the allocation domain since the scheduler takes charge of setting the allocation domain in the registers. Based on the difference in the (maybe)allocation domain of the producer and consumer of the second LoadStoreOp, we may do transposed load when reading to registers.  See also #2315.
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
4 participants
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.
This PR modifies how we determine when we do a LdMatrix vs LdMatrixTranspose.
We look at the (maybe) allocation domain of the producer and consumer of the load store op do determine if a transpose is required. We get the allocation domain of the consumer and check if the innermost dimension maps to the innermost dimension of the allocation domain of the producer. We transpose if it doesn't map.