Skip to content
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

Merged

Conversation

zxybazh
Copy link
Member

@zxybazh zxybazh commented Mar 20, 2023

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

@tvm-bot
Copy link
Collaborator

tvm-bot commented Mar 20, 2023

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

Copy link
Contributor

@sunggg sunggg left a 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(
Copy link
Contributor

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?

Copy link
Member Author

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.

@zxybazh zxybazh force-pushed the bugfix/2023-03-20/fix-data-type-lower-tvm-builtin branch from 2ad2de4 to 1d335ec Compare March 21, 2023 00:16

Verified

This commit was signed with the committer’s verified signature.
pradyunsg Pradyun Gedam
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
Copy link
Contributor

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.

Copy link
Member Author

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.

@zxybazh zxybazh force-pushed the bugfix/2023-03-20/fix-data-type-lower-tvm-builtin branch from 1d335ec to 1ee7cb1 Compare March 21, 2023 00:21
@zxybazh zxybazh merged commit 7f6da09 into apache:main Mar 21, 2023
vinx13 pushed a commit to vinx13/tvm that referenced this pull request Mar 27, 2023
Fix data type and add minimal reproducible test.

Co-authored-by: Sunghyun Park <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

6 participants