You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "Cookiee235 (via GitHub)" <gi...@apache.org> on 2024/09/04 16:54:24 UTC
[I] [Bug] inferencing on the CUDA and LLVM produce inconsistent results [tvm]
Cookiee235 opened a new issue, #17333:
URL: https://github.com/apache/tvm/issues/17333
### Actual behavior
```
Traceback (most recent call last):
File "/share_container/optfuzz/res/bugs/inconsis333.py", line 266, in <module>
np.testing.assert_allclose(outputs1, outputs2,1e-3,1e-3)
File "/root/miniconda3/lib/python3.12/site-packages/numpy/testing/_private/utils.py", line 1504, in assert_allclose
assert_array_compare(compare, actual, desired, err_msg=str(err_msg),
File "/root/miniconda3/lib/python3.12/contextlib.py", line 81, in inner
return func(*args, **kwds)
^^^^^^^^^^^^^^^^^^^
File "/root/miniconda3/lib/python3.12/site-packages/numpy/testing/_private/utils.py", line 797, in assert_array_compare
raise AssertionError(msg)
AssertionError:
Not equal to tolerance rtol=0.001, atol=0.001
Mismatched elements: 169 / 200 (84.5%)
Max absolute difference: 2013265920
Max relative difference: -0.
x: array([[ 2, 8, 24, 64, 160,
384, 896, 2048, 4608, 10240,
22528, 49152, 106496, 229376, 491520,...
y: array([[ 2, 8, 24, 64, 160,
384, 896, 2048, 4608, 10240,
22528, 49152, 106496, 229376, 491520,...
```
### Steps to reproduce
This is a complex Relax IR, but I cannot simplify it further. Further reducing the Relax IR will lead to the disappearance of the inconsistency
<details>
<summary>Reproducible script</summary>
```
import tvm
from tvm import relax
import numpy as np
import tvm
metadata = tvm.ir.load_json("""{
\"root\": 1,
\"nodes\": [
{
\"type_key\": \"\"
},
{
\"type_key\": \"Map\",
\"keys\": [
\"relax.expr.Constant\"
],
\"data\": [2]
},
{
\"type_key\": \"Array\",
\"data\": [3]
},
{
\"type_key\": \"relax.expr.Constant\",
\"attrs\": {
\"_checked_type_\": \"11\",
\"data\": \"0\",
\"span\": \"0\",
\"struct_info_\": \"4\"
}
},
{
\"type_key\": \"relax.TensorStructInfo\",
\"attrs\": {
\"dtype\": \"float32\",
\"ndim\": \"2\",
\"shape\": \"5\",
\"span\": \"0\",
\"vdevice\": \"0\"
}
},
{
\"type_key\": \"relax.expr.ShapeExpr\",
\"attrs\": {
\"_checked_type_\": \"10\",
\"span\": \"0\",
\"struct_info_\": \"9\",
\"values\": \"6\"
}
},
{
\"type_key\": \"Array\",
\"data\": [7, 8]
},
{
\"type_key\": \"IntImm\",
\"attrs\": {
\"dtype\": \"int64\",
\"span\": \"0\",
\"value\": \"16\"
}
},
{
\"type_key\": \"IntImm\",
\"attrs\": {
\"dtype\": \"int64\",
\"span\": \"0\",
\"value\": \"16\"
}
},
{
\"type_key\": \"relax.ShapeStructInfo\",
\"attrs\": {
\"ndim\": \"2\",
\"span\": \"0\",
\"values\": \"6\"
}
},
{
\"type_key\": \"relax.ShapeType\",
\"attrs\": {
\"ndim\": \"2\",
\"span\": \"0\"
}
},
{
\"type_key\": \"relax.DynTensorType\",
\"attrs\": {
\"dtype\": \"float32\",
\"ndim\": \"2\",
\"span\": \"0\"
}
}
],
\"b64ndarrays\": [
\"P6G0lvBAXt0AAAAAAAAAAAEAAAAAAAAAAgAAAAIgAQAQAAAAAAAAABAAAAAAAAAAAAQAAAAAAAAAAAAAAACAPwAAAEAAAEBAAACAQAAAoEAAAMBAAADgQAAAAEEAABBBAAAgQQAAMEEAAEBBAABQQQAAYEEAAHBBAACAQQAAiEEAAJBBAACYQQAAoEEAAKhBAACwQQAAuEEAAMBBAADIQQAA0EEAANhBAADgQQAA6EEAAPBBAAD4QQAAAEIAAARCAAAIQgAADEIAABBCAAAUQgAAGEIAABxCAAAgQgAAJEIAAChCAAAsQgAAMEIAADRCAAA4QgAAPEIAAEBCAABEQgAASEIAAExCAABQQgAAVEIAAFhCAABcQgAAYEIAAGRCAABoQgAAbEIAAHBCAAB0QgAAeEIAAHxCAACAQgAAgkIAAIRCAACGQgAAiEIAAIpCAACMQgAAjkIAAJBCAACSQgAAlEIAAJZCAACYQgAAmkIAAJxCAACeQgAAoEIAAKJCAACkQgAApkIAAKhCAACqQgAArEIAAK5CAACwQgAAskIAALRCAAC2QgAAuEIAALpCAAC8QgAAvkIAAMBCAADCQgAAxEIAAMZCAADIQgAAykIAAMxCAADOQgAA0EIAANJCAADUQgAA1kIAANhCAADaQgAA3EIAAN5CAADgQgAA4kIAAORCAADmQgAA6EIAAOpCAADsQgAA7kIAAPBCAADyQgAA9EIAAPZCAAD4QgAA+kIAAPxCAAD+QgAAAEMAAAFDAAACQwAAA0MAAARDAAAFQwAABkMAAAdDAAAIQwAACUMAAApDAAALQwAADEMAAA1DAAAOQwAAD0MAABBDAAARQwAAEkMAABNDAAAUQwAAFUMAABZDAAAXQwAAGEMAABlDAAAaQwAAG0MAABxDAAAdQwAAHkMAAB9DAAAgQwAAIUMAACJDAAAjQwAAJEMAACVDAAAmQwAAJ0MAAChDAAApQwAAKkMAA
CtDAAAsQwAALUMAAC5DAAAvQwAAMEMAADFDAAAyQwAAM0MAADRDAAA1QwAANkMAADdDAAA4QwAAOUMAADpDAAA7QwAAPEMAAD1DAAA+QwAAP0MAAEBDAABBQwAAQkMAAENDAABEQwAARUMAAEZDAABHQwAASEMAAElDAABKQwAAS0MAAExDAABNQwAATkMAAE9DAABQQwAAUUMAAFJDAABTQwAAVEMAAFVDAABWQwAAV0MAAFhDAABZQwAAWkMAAFtDAABcQwAAXUMAAF5DAABfQwAAYEMAAGFDAABiQwAAY0MAAGRDAABlQwAAZkMAAGdDAABoQwAAaUMAAGpDAABrQwAAbEMAAG1DAABuQwAAb0MAAHBDAABxQwAAckMAAHNDAAB0QwAAdUMAAHZDAAB3QwAAeEMAAHlDAAB6QwAAe0MAAHxDAAB9QwAAfkMAAH9D\"
],
\"attrs\": {\"tvm_version\": \"0.18.dev0\"}
}""")
from tvm.script import ir as I
from tvm.script import tir as T
from tvm.script import relax as R
@I.ir_module
class Module:
@T.prim_func(private=True)
def add(x: T.Buffer((T.int64(10), T.int64(20)), "int32"), B: T.Buffer((), "int32"), T_add: T.Buffer((T.int64(10), T.int64(20)), "int32")):
T.func_attr({"op_pattern": 0, "tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0, ax1 in T.grid(T.int64(10), T.int64(20)):
with T.block("T_add"):
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(x[v_ax0, v_ax1], B[()])
T.writes(T_add[v_ax0, v_ax1])
T_add[v_ax0, v_ax1] = x[v_ax0, v_ax1] + B[()]
@T.prim_func(private=True)
def cast(res: T.Buffer((T.int64(16), T.int64(16)), "float32"), compute: T.Buffer((T.int64(16), T.int64(16)), "int32")):
T.func_attr({"tir.noalias": T.bool(True)})
# with T.block("root"):
for i0, i1 in T.grid(T.int64(16), T.int64(16)):
with T.block("compute"):
v_i0, v_i1 = T.axis.remap("SS", [i0, i1])
T.reads(res[v_i0, v_i1])
T.writes(compute[v_i0, v_i1])
compute[v_i0, v_i1] = T.Cast("int32", res[v_i0, v_i1])
@T.prim_func
def identity(A: T.Buffer((16, 16), "float32"), B: T.Buffer((16, 16), "float32")):
# with T.block("root"):
for i, j in T.grid(16, 16):
with T.block("identity"):
vi, vj = T.axis.remap("SS", [i, j])
T.reads(A[vi, vj])
T.writes(B[vi, vj])
B[vi, vj] = A[vi, vj]
@T.prim_func(private=True)
def left_shift(lv1: T.Buffer((T.int64(10), T.int64(20)), "int32"), lv3: T.Buffer((T.int64(10), T.int64(20)), "int32"), T_left_shift: T.Buffer((T.int64(10), T.int64(20)), "int32")):
T.func_attr({"op_pattern": 0, "tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0, ax1 in T.grid(T.int64(10), T.int64(20)):
with T.block("T_left_shift"):
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(lv1[v_ax0, v_ax1], lv3[v_ax0, v_ax1])
T.writes(T_left_shift[v_ax0, v_ax1])
T_left_shift[v_ax0, v_ax1] = T.shift_left(lv1[v_ax0, v_ax1], lv3[v_ax0, v_ax1])
@T.prim_func(private=True)
def reshape(res: T.Buffer((T.int64(16), T.int64(16)), "int32"), T_reshape: T.Buffer((T.int64(256),), "int32")):
T.func_attr({"tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0 in range(T.int64(256)):
with T.block("T_reshape"):
v_ax0 = T.axis.spatial(T.int64(256), ax0)
T.reads(res[v_ax0 % T.int64(256) // T.int64(16), v_ax0 % T.int64(16)])
T.writes(T_reshape[v_ax0])
T_reshape[v_ax0] = res[v_ax0 % T.int64(256) // T.int64(16), v_ax0 % T.int64(16)]
@T.prim_func(private=True)
def reshape1(temp: T.Buffer((T.int64(200),), "int32"), T_reshape: T.Buffer((T.int64(10), T.int64(20)), "int32")):
T.func_attr({"tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0, ax1 in T.grid(T.int64(10), T.int64(20)):
with T.block("T_reshape"):
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(temp[(v_ax0 * T.int64(20) + v_ax1) % T.int64(200)])
T.writes(T_reshape[v_ax0, v_ax1])
T_reshape[v_ax0, v_ax1] = temp[(v_ax0 * T.int64(20) + v_ax1) % T.int64(200)]
@T.prim_func(private=True)
def squeeze(lv: T.Buffer((T.int64(10), T.int64(20)), "int32"), T_squeeze: T.Buffer((T.int64(10), T.int64(20)), "int32")):
T.func_attr({"op_pattern": 0, "tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0, ax1 in T.grid(T.int64(10), T.int64(20)):
with T.block("T_squeeze"):
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(lv[v_ax0, v_ax1])
T.writes(T_squeeze[v_ax0, v_ax1])
T_squeeze[v_ax0, v_ax1] = lv[v_ax0, v_ax1]
@T.prim_func(private=True)
def strided_slice(tensor_1dim: T.Buffer((T.int64(256),), "int32"), T_strided_slice_with_axes: T.Buffer((T.int64(200),), "int32")):
T.func_attr({"tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0 in range(T.int64(200)):
with T.block("T_strided_slice_with_axes"):
v_ax0 = T.axis.spatial(T.int64(200), ax0)
T.reads(tensor_1dim[v_ax0])
T.writes(T_strided_slice_with_axes[v_ax0])
T_strided_slice_with_axes[v_ax0] = tensor_1dim[v_ax0]
@T.prim_func(private=True)
def transpose(lv: T.Buffer((T.int64(10), T.int64(20)), "int32"), T_transpose: T.Buffer((T.int64(20), T.int64(10)), "int32")):
T.func_attr({"op_pattern": 2, "tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0, ax1 in T.grid(T.int64(20), T.int64(10)):
with T.block("T_transpose"):
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(lv[v_ax1, v_ax0])
T.writes(T_transpose[v_ax0, v_ax1])
T_transpose[v_ax0, v_ax1] = lv[v_ax1, v_ax0]
@T.prim_func(private=True)
def transpose1(lv2: T.Buffer((T.int64(20), T.int64(10)), "int32"), T_transpose: T.Buffer((T.int64(10), T.int64(20)), "int32")):
T.func_attr({"op_pattern": 2, "tir.noalias": T.bool(True)})
# with T.block("root"):
for ax0, ax1 in T.grid(T.int64(10), T.int64(20)):
with T.block("T_transpose"):
v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
T.reads(lv2[v_ax1, v_ax0])
T.writes(T_transpose[v_ax0, v_ax1])
T_transpose[v_ax0, v_ax1] = lv2[v_ax1, v_ax0]
@R.function
def foo(x: R.Tensor((10, 20), dtype="int32")) -> R.Tensor((10, 20), dtype="int32"):
cls = Module
with R.dataflow():
lv = R.call_tir(cls.add, (x, R.const(1, "int32")), out_sinfo=R.Tensor((10, 20), dtype="int32"))
lv1 = R.call_tir(cls.squeeze, (lv,), out_sinfo=R.Tensor((10, 20), dtype="int32"))
lv2 = R.call_tir(cls.transpose, (lv,), out_sinfo=R.Tensor((20, 10), dtype="int32"))
lv3 = R.call_tir(cls.transpose1, (lv2,), out_sinfo=R.Tensor((10, 20), dtype="int32"))
gv = R.call_tir(cls.left_shift, (lv1, lv3), out_sinfo=R.Tensor((10, 20), dtype="int32"))
R.output(gv)
return gv
@R.function
def func1() -> R.Tensor((16, 16), dtype="float32"):
cls = Module
with R.dataflow():
gv0 = R.call_tir(cls.identity, (metadata["relax.expr.Constant"][0],), out_sinfo=R.Tensor((16, 16), dtype="float32"))
R.output(gv0)
return gv0
@R.function
def main(A: R.Tensor((16,16), dtype="float32")) -> R.Tensor((10, 20), dtype="int32"):
cls = Module
res: R.Tensor((16, 16), dtype="float32") = cls.func1()
res_1 = R.call_tir(cls.cast, (res,), out_sinfo=R.Tensor((16, 16), dtype="int32"))
tensor_1dim = R.call_tir(cls.reshape, (res_1,), out_sinfo=R.Tensor((256,), dtype="int32"))
temp = R.call_tir(cls.strided_slice, (tensor_1dim,), out_sinfo=R.Tensor((200,), dtype="int32"))
para0 = R.call_tir(cls.reshape1, (temp,), out_sinfo=R.Tensor((10, 20), dtype="int32"))
res_2: R.Tensor((10, 20), dtype="int32") = cls.foo(para0)
return res_2
mod = Module
mod = relax.transform.FuseTIR()(mod)
def compile_mod(mod, func_name, target, *inputs):
if target == 'llvm':
ex = relax.build(mod, target='llvm')
vm = relax.VirtualMachine(ex, tvm.cpu())
else:
with tvm.target.Target("cuda"):
mod = tvm.tir.transform.DefaultGPUSchedule()(mod)
ex = relax.build(mod, target='cuda')
vm = relax.VirtualMachine(ex, tvm.gpu())
mod_outputs = vm[f'{func_name}'](*inputs)
mod_outputs = mod_outputs.numpy()
return mod_outputs
input_0 = np.random.random([16,16]).astype('float32')
outputs1 = compile_mod(mod, 'main', 'llvm', tvm.nd.array(input_0))
print(outputs1)
outputs2 = compile_mod(mod, 'main', 'cuda', tvm.nd.array(input_0,tvm.cuda()))
print(outputs2)
np.testing.assert_allclose(outputs1, outputs2,1e-3,1e-3)
```
</details>
@Lunderberg @junrushao
--
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.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [I] [Bug] inferencing on the CUDA and LLVM produce inconsistent results [tvm]
Posted by "Cookiee235 (via GitHub)" <gi...@apache.org>.
Cookiee235 closed issue #17333: [Bug] inferencing on the CUDA and LLVM produce inconsistent results
URL: https://github.com/apache/tvm/issues/17333
--
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
Re: [I] [Bug] inferencing on the CUDA and LLVM produce inconsistent results [tvm]
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on issue #17333:
URL: https://github.com/apache/tvm/issues/17333#issuecomment-2329813598
Looks like the tensor specifying the number of bits to shift was populated with values from 0 to 255. There's some undefined behavior in C's definition of `a << b`, which is inherited by LLVM, where the output is undefined when the number of bits to shift is greater than or equal to the size of the value. Here, the first 31 values are matched between the two outputs, but values after that, produced by left-shift that is larger than an `int32`, runs into the undefined behavior.
--
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
Re: [I] [Bug] inferencing on the CUDA and LLVM produce inconsistent results [tvm]
Posted by "Cookiee235 (via GitHub)" <gi...@apache.org>.
Cookiee235 commented on issue #17333:
URL: https://github.com/apache/tvm/issues/17333#issuecomment-2331759003
@Lunderberg Thanks for your deep investigation. I closed it due to the undefined behavior. Thanks again!
--
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