The subtraction problem when using the unsigned integer

Questions
#1

Hello, I’m studying TVM and its low-level operators recently. These days I have tried to construct some test files and I found the returned value maybe not accurate when I used the unsigned integer in the tir.sub python api. For example, when I used the unsigned 32 bit value 3 minus unsigned 32 bit value 5, the TVM returned 4294967294.It seems like the buffer overflow problem. This question occurred in the TVM 0.11dev0, the git commit is 5019dcee0efb526cc26858eb87e38f32e4f1a464. Should such a behavior be considered a bug? Thanks.

import tvm
from tvm import tir,te

f = tir.const(3,'uint32')
g = tir.const(5,'uint32')
output = tir.ret(f-g)
output = tir.Evaluate(output)
func = tir.PrimFunc([],output)
mod = tvm.build(func)
print(mod())
#2

When I tested this file in other versions, for example TVM 0.10 and TVM 0.9 version, TVM threw the check failed error messages when I executed the test file. Probably this problem is caused by the code update. In the TVM 0.11dev0 version or the latest TVM version, the subtraction api uses the GetFoldResultInt64Repr function when the data type is uint32, but GetFoldResultInt64Repr function seems not check whether the unsigned integer value is negative.

inline int64_t GetFoldResultInt64Repr(int64_t x, const DataType& dtype) {
  if (dtype.bits() < 64) {
    x &= (1LL << dtype.bits()) - 1;
  }
  if (dtype.is_int()) {
    int64_t m = 1LL << (dtype.bits() - 1);
    x = (x ^ m) - m;
  }
  return x;
}
#3

Thank you for your findings and your reporting.

I think the behavior of 0.11 makes more sense to me as the native cpp code will not check the negative unsigned integer :slight_smile: It should be the user’s/ developer’s responsibility when using unsigned integers.

#4

I agree that it should be the user’s responsibility and an error is probably too restrictive, but perhaps in case of constants, we can show some warnings as overflows are rarely what a user intends.

It might help make it easier to debug correctness issues.

#5

@Hzfengsy @sanirudh Thanks for your reply! :grinning:

#6

Here is another buffer overflow example and this error has been fixed in the latest version, maybe we should take a uniform standard and add some strict checking? :slightly_smiling_face:

import tvm
from tvm import tir
a=tvm.tir.const(-1,'uint32')
out=tir.ret(a)
out=tir.Evaluate(out)
func=tir.PrimFunc({},out)
mod=tvm.build(func)
print(mod())