-
Notifications
You must be signed in to change notification settings - Fork 3.5k
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
[TIR] Fix Datatype in Lower TVM Builtin #14347
[TIR] Fix Datatype in Lower TVM Builtin #14347
Conversation
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.
Generated by tvm-bot |
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.
Thanks, @zxybazh!
Just quick comment.
@@ -172,6 +172,47 @@ def build_tir(): | |||
tvm.testing.assert_allclose(a.numpy(), expected_value) | |||
|
|||
|
|||
def test_lower_overflow_int32(): | |||
@T.prim_func | |||
def variance4( |
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.
Can we only keep the minimal reproducible parts?
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.
Thanks, have updated to the minimal version.
2ad2de4
to
1d335ec
Compare
Co-authored-by: Sunghyun Park <[email protected]>
for (size_t i = 0; i < op->extents.size(); ++i) { | ||
// set total_bytes to uint64 to avoid overflow |
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.
possible for us to also add a guard check here to verify that total_bytes isn't negative? although...i guess it is kinda unlikely to alloc > (1 << 31) worth of space.
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.
I think technically it's possible to allocate 2G of memory and also possible to have an intermediate large result between passes. On the other hand each loop extent are supposed to be positive so unless overflow it's unlikely to produce any negative results. I also tried locally to multiply uint64 with negative int32, turns out it will trigger an error. Therefore I think we are good with current change here.
1d335ec
to
1ee7cb1
Compare
Fix data type and add minimal reproducible test. Co-authored-by: Sunghyun Park <[email protected]>
This PR fixes the issue of overflow when calculating the size of buffer using product of loop extents. In certain cases it could overflow to a negative number and fail to compile. Unit test is also attached.
Co-authored with Sunghyun Park <[email protected]>
CC: @sunggg @jwfromm @tqchen