You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "zxybazh (via GitHub)" <gi...@apache.org> on 2023/03/20 23:54:12 UTC

[GitHub] [tvm] zxybazh opened a new pull request, #14347: [TIR] Fix Datatype in Lower TVM Builtin

zxybazh opened a new pull request, #14347:
URL: https://github.com/apache/tvm/pull/14347

   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.
   
   CC: @sunggg @jwfromm @tqchen 


-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


[GitHub] [tvm] sunggg commented on a diff in pull request #14347: [TIR] Fix Datatype in Lower TVM Builtin

Posted by "sunggg (via GitHub)" <gi...@apache.org>.
sunggg commented on code in PR #14347:
URL: https://github.com/apache/tvm/pull/14347#discussion_r1142781944


##########
tests/python/unittest/test_tir_transform_lower_tvm_builtin.py:
##########
@@ -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(

Review Comment:
   Can we only keep the minimal reproducible parts? 



-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


[GitHub] [tvm] zxybazh commented on a diff in pull request #14347: [TIR] Fix Datatype in Lower TVM Builtin

Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14347:
URL: https://github.com/apache/tvm/pull/14347#discussion_r1142786040


##########
tests/python/unittest/test_tir_transform_lower_tvm_builtin.py:
##########
@@ -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(

Review Comment:
   Thanks, have updated to the minimal version.



-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


[GitHub] [tvm] zxybazh merged pull request #14347: [TIR] Fix Datatype in Lower TVM Builtin

Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh merged PR #14347:
URL: https://github.com/apache/tvm/pull/14347


-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


[GitHub] [tvm] zxybazh commented on a diff in pull request #14347: [TIR] Fix Datatype in Lower TVM Builtin

Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on code in PR #14347:
URL: https://github.com/apache/tvm/pull/14347#discussion_r1142799742


##########
src/tir/transforms/lower_tvm_builtin.cc:
##########
@@ -239,8 +239,9 @@ class BuiltinLower : public StmtExprMutator {
         }
       }
     }
-    PrimExpr total_bytes = make_const(op->extents[0].dtype(), nbytes);
+    PrimExpr total_bytes = make_const(DataType::UInt(64), nbytes);
     for (size_t i = 0; i < op->extents.size(); ++i) {
+      // set total_bytes to uint64 to avoid overflow

Review Comment:
   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.



-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


[GitHub] [tvm] areusch commented on a diff in pull request #14347: [TIR] Fix Datatype in Lower TVM Builtin

Posted by "areusch (via GitHub)" <gi...@apache.org>.
areusch commented on code in PR #14347:
URL: https://github.com/apache/tvm/pull/14347#discussion_r1142791245


##########
src/tir/transforms/lower_tvm_builtin.cc:
##########
@@ -239,8 +239,9 @@ class BuiltinLower : public StmtExprMutator {
         }
       }
     }
-    PrimExpr total_bytes = make_const(op->extents[0].dtype(), nbytes);
+    PrimExpr total_bytes = make_const(DataType::UInt(64), nbytes);
     for (size_t i = 0; i < op->extents.size(); ++i) {
+      // set total_bytes to uint64 to avoid overflow

Review Comment:
   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.



-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


[GitHub] [tvm] tvm-bot commented on pull request #14347: [TIR] Fix Datatype in Lower TVM Builtin

Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #14347:
URL: https://github.com/apache/tvm/pull/14347#issuecomment-1477098742

   <!---bot-comment-->
   
   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](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
   
   <!--bot-comment-ccs-start-->
    * cc @Hzfengsy, @junrushao, @quic-sanirudh, @shingjan <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
   
   <sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>


-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org