-
Notifications
You must be signed in to change notification settings - Fork 92
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[PTX] Enable migration of ldmatrix #2692
base: SYCLomatic
Are you sure you want to change the base?
Conversation
bd88679
to
4492ec9
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
@@ -557,12 +557,15 @@ bool SYCLGenBase::emitVectorType(const InlineAsmVectorType *T) { | |||
OS() << ", "; | |||
switch (T->getKind()) { | |||
case InlineAsmVectorType::v2: | |||
case InlineAsmVectorType::x1: |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
What's the relation between ptx x1
x2
x4
and vector size? And should have a lit test for it.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
In "xn", n refers to the number of output args being passed to the instruction
There seems to be no lit test added for "vn", so I'm not sure about lit test for this scenario
What do you advice?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
By the way, I don't see the usage of the vector size, is that right?
OS() << ", "; | ||
if (const auto *VE = | ||
dyn_cast<InlineAsmVectorExpr>(Inst->getOutputOperand())) { | ||
for (unsigned Inst = 0; Inst != VE->getNumElements(); ++Inst) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Should we check the vector size satisfy the instruction requirement?
This PR adds support for the migration of ldmatrix PTX ASM API