-
Notifications
You must be signed in to change notification settings - Fork 54
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
Fix alias analysis #185
Fix alias analysis #185
Conversation
Set was replaced with LoadStoreOp, which was not recognized as UnaryOp, and thus the alias analysis failed to detect safe aliasing
csrc/ir_utils.cpp
Outdated
@@ -572,6 +572,11 @@ bool isReductionTvOp(const Expr* expr) { | |||
return ir_utils::isTvOp(expr) && isReductionOp(expr); | |||
} | |||
|
|||
bool isPointwiseTvOp(const Expr* expr) { | |||
return isTvOp(expr) && | |||
expr->isOneOf<UnaryOp, BinaryOp, TernaryOp, LoadStoreOp>(); |
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.
Just raise a question for discussion: Should we disable inner sharing for LoadStoreOp
with consumer that has rfactor domain (that is, permutation ops)? I am not sure, but I think it should be OK to share.
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.
T2 = transpose(T1)
T3 = sin(T2)
For the example above, I think it is OK to share T3 with T1.
We shouldn't share T2
with T1
, but I think we are already disallowing this type of sharing.
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.
That's a good point. Is it really safe? Suppose t1 = transpose(t0)
, and both of them are on shared memory, code like below seemed to be generated:
__shared__ t0[N] = ...;
auto& t1 = t0;
t1[i * k + j] = t0[j *k + i];
This isn't valid an it's not really pointwise, is 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.
Wait what, is t1 sharing with t0? I think we previously disabled it. csarofeen/pytorch#2490
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.
That is right. But should we call it a pointwise op?
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.
Probably not
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.
Updated.
!build |
!build |
Set was replaced with LoadStoreOp, which was not recognized as UnaryOp, and thus the alias analysis failed to detect safe aliasing
Fixes #163