talha-ahsan opened a new issue, #17386:
URL: https://github.com/apache/tvm/issues/17386
### Expected behavior
A successful compilation, or messages about what part of the program being
compiled fails. As far as I can tell this _should_ at least surface a more
user-friendly message if the TIR is invalid.
### Actual behavior
I get a long trace (added below) which ends with this message:
`InternalError: Check failed: (is_zero(op->min)) is false:`
This seems to fail in the `tir.LowerIntrin` pass
### Steps to reproduce
This script causes the exception to occur:
```python
import tvm
from tvm import tir
from tvm.tir.analysis.analysis import verify_well_formed, verify_memory
from tvm.script import tir as T
@T.prim_func
def tvmgen_default_fused_nn_conv2d_3(p0: T.Buffer((1, 256, 56, 56),
"float32"), p1: T.Buffer((256, 256, 3, 3), "float32"), output_unpack:
T.Buffer((1, 256, 56, 56), "float32")):
T.func_attr({"from_legacy_te_schedule": T.bool(True), "hash":
"032dbe00302af996", "target": T.target({"host": {"keys": ["cpu"], "kind":
"llvm", "tag": ""}, "keys": ["cpu"], "kind": "llvm", "tag": ""}),
"tir.noalias": T.bool(True)})
data_vec = T.allocate([802816], "float32", "global")
data_pad = T.allocate([861184], "float32", "global")
data_vec_1 = T.Buffer((802816,), data=data_vec)
for bs_c_fused_h_fused in T.parallel(3584):
for w, vc in T.grid(56, 4):
p0_1 = T.Buffer((802816,), data=p0.data)
data_vec_1[bs_c_fused_h_fused * 224 + w * 4 + vc] =
p0_1[bs_c_fused_h_fused // 56 * 12544 + vc * 3136 + bs_c_fused_h_fused % 56 *
56 + w]
data_pad_1 = T.Buffer((861184,), data=data_pad)
for i0_i1_fused_i2_fused in T.parallel(3712):
for i3 in range(58):
cse_var_2: T.int32 = i0_i1_fused_i2_fused % 58
cse_var_1: T.int32 = i3 * 4
data_pad_1[i0_i1_fused_i2_fused * 232 +
cse_var_1:i0_i1_fused_i2_fused * 232 + cse_var_1 + 4] = T.if_then_else(1 <=
cse_var_2 and cse_var_2 < 57 and 1 <= i3 and i3 < 57,
data_vec_1[i0_i1_fused_i2_fused // 58 * 12544 + cse_var_2 * 224 + cse_var_1 -
228:i0_i1_fused_i2_fused // 58 * 12544 + cse_var_2 * 224 + cse_var_1 - 228 +
4], T.Broadcast(T.float32(0), 4))
data_vec_2 = T.Buffer((589824,), data=data_vec)
for occ_k_h_fused in T.parallel(192):
for icc, k_w, icb in T.grid(64, 3, 4):
cse_var_4: T.int32 = occ_k_h_fused % 3
cse_var_3: T.int32 = occ_k_h_fused // 3 * 9216
p1_1 = T.Buffer((589824,), data=p1.data)
data_vec_2[cse_var_3 + icc * 144 + cse_var_4 * 48 + k_w * 16 +
icb * 4:cse_var_3 + icc * 144 + cse_var_4 * 48 + k_w * 16 + icb * 4 + 4] =
p1_1[cse_var_3 + icc * 36 + icb * 9 + cse_var_4 * 3 + k_w:cse_var_3 + icc * 36
+ icb * 9 + cse_var_4 * 3 + k_w + 9216:2304]
for n_c_outer_fused_h_fused in T.parallel(3584):
conv2d_NCHWc = T.allocate([56], "float32x4", "global")
conv2d_NCHWc_global = T.allocate([28], "float32x4", "global")
conv2d_NCHWc_1 = T.Buffer((56,), "float32x4", data=conv2d_NCHWc)
for ow_outer in range(2):
conv2d_NCHWc_global_1 = T.Buffer((28,), "float32x4",
data=conv2d_NCHWc_global)
conv2d_NCHWc_global_1[0] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[1] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[2] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[3] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[4] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[5] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[6] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[7] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[8] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[9] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[10] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[11] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[12] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[13] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[14] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[15] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[16] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[17] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[18] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[19] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[20] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[21] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[22] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[23] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[24] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[25] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[26] = T.Broadcast(T.float32(0), 4)
conv2d_NCHWc_global_1[27] = T.Broadcast(T.float32(0), 4)
for ic_outer, kh, kw, ic_inner in T.grid(64, 3, 3, 4):
cse_var_6: T.int32 = n_c_outer_fused_h_fused // 56 * 9216 +
ic_outer * 144 + kh * 48 + kw * 16 + ic_inner * 4
cse_var_5: T.int32 = ic_outer * 13456 + kh * 232 +
n_c_outer_fused_h_fused % 56 * 232 + ow_outer * 112 + kw * 4 + ic_inner
conv2d_NCHWc_global_1[0] = conv2d_NCHWc_global_1[0] +
T.Broadcast(data_pad_1[cse_var_5], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[1] = conv2d_NCHWc_global_1[1] +
T.Broadcast(data_pad_1[cse_var_5 + 4], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[2] = conv2d_NCHWc_global_1[2] +
T.Broadcast(data_pad_1[cse_var_5 + 8], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[3] = conv2d_NCHWc_global_1[3] +
T.Broadcast(data_pad_1[cse_var_5 + 12], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[4] = conv2d_NCHWc_global_1[4] +
T.Broadcast(data_pad_1[cse_var_5 + 16], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[5] = conv2d_NCHWc_global_1[5] +
T.Broadcast(data_pad_1[cse_var_5 + 20], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[6] = conv2d_NCHWc_global_1[6] +
T.Broadcast(data_pad_1[cse_var_5 + 24], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[7] = conv2d_NCHWc_global_1[7] +
T.Broadcast(data_pad_1[cse_var_5 + 28], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[8] = conv2d_NCHWc_global_1[8] +
T.Broadcast(data_pad_1[cse_var_5 + 32], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[9] = conv2d_NCHWc_global_1[9] +
T.Broadcast(data_pad_1[cse_var_5 + 36], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[10] = conv2d_NCHWc_global_1[10] +
T.Broadcast(data_pad_1[cse_var_5 + 40], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[11] = conv2d_NCHWc_global_1[11] +
T.Broadcast(data_pad_1[cse_var_5 + 44], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[12] = conv2d_NCHWc_global_1[12] +
T.Broadcast(data_pad_1[cse_var_5 + 48], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[13] = conv2d_NCHWc_global_1[13] +
T.Broadcast(data_pad_1[cse_var_5 + 52], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[14] = conv2d_NCHWc_global_1[14] +
T.Broadcast(data_pad_1[cse_var_5 + 56], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[15] = conv2d_NCHWc_global_1[15] +
T.Broadcast(data_pad_1[cse_var_5 + 60], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[16] = conv2d_NCHWc_global_1[16] +
T.Broadcast(data_pad_1[cse_var_5 + 64], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[17] = conv2d_NCHWc_global_1[17] +
T.Broadcast(data_pad_1[cse_var_5 + 68], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[18] = conv2d_NCHWc_global_1[18] +
T.Broadcast(data_pad_1[cse_var_5 + 72], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[19] = conv2d_NCHWc_global_1[19] +
T.Broadcast(data_pad_1[cse_var_5 + 76], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[20] = conv2d_NCHWc_global_1[20] +
T.Broadcast(data_pad_1[cse_var_5 + 80], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[21] = conv2d_NCHWc_global_1[21] +
T.Broadcast(data_pad_1[cse_var_5 + 84], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[22] = conv2d_NCHWc_global_1[22] +
T.Broadcast(data_pad_1[cse_var_5 + 88], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[23] = conv2d_NCHWc_global_1[23] +
T.Broadcast(data_pad_1[cse_var_5 + 92], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[24] = conv2d_NCHWc_global_1[24] +
T.Broadcast(data_pad_1[cse_var_5 + 96], 4) * data_vec_2[cse_var_6:cse_var_6 + 4]
conv2d_NCHWc_global_1[25] = conv2d_NCHWc_global_1[25] +
T.Broadcast(data_pad_1[cse_var_5 + 100], 4) * data_vec_2[cse_var_6:cse_var_6 +
4]
conv2d_NCHWc_global_1[26] = conv2d_NCHWc_global_1[26] +
T.Broadcast(data_pad_1[cse_var_5 + 104], 4) * data_vec_2[cse_var_6:cse_var_6 +
4]
conv2d_NCHWc_global_1[27] = conv2d_NCHWc_global_1[27] +
T.Broadcast(data_pad_1[cse_var_5 + 108], 4) * data_vec_2[cse_var_6:cse_var_6 +
4]
for ow_inner in range(28):
conv2d_NCHWc_1[ow_outer * 28 + ow_inner] =
conv2d_NCHWc_global_1[ow_inner]
for w_outer in range(2):
for w_inner in range(-1075824214, -1075824186):
cse_var_7: T.int32 = w_outer * 28
output_unpack_1 = T.Buffer((802816,),
data=output_unpack.data)
output_unpack_1[n_c_outer_fused_h_fused // 56 * 12544 +
n_c_outer_fused_h_fused % 56 * 56 + cse_var_7 + w_inner:n_c_outer_fused_h_fused
// 56 * 12544 + n_c_outer_fused_h_fused % 56 * 56 + cse_var_7 + w_inner +
12544:3136] = conv2d_NCHWc_1[cse_var_7 + cse_var_7]
func = tvmgen_default_fused_nn_conv2d_3
mod = tvm.ir.IRModule({'main': func})
if not verify_well_formed(mod) and verify_memory(func):
print("Validation failed")
else:
with tvm.transform.PassContext(opt_level=0):
nopt_mod = tvm.build(mod)
print("Success!")
```
### Dump
```
Traceback (most recent call last):
File "<path_to_script>/TVMBugReport3/reprod.py", line 113, in <module>
nopt_mod = tvm.build(mod)
File "<path_to_tvm>/tvm/python/tvm/driver/build_module.py", line 297, in
build
rt_mod_host = _driver_ffi.tir_to_runtime(annotated_mods, target_host)
File "<path_to_tvm>/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line 239,
in __call__
raise_last_ffi_error()
File "<path_to_tvm>/tvm/python/tvm/_ffi/base.py", line 481, in
raise_last_ffi_error
raise py_err
tvm.error.InternalError: Traceback (most recent call last):
119:
tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::runtime::Module
(tvm::runtime::Map<tvm::Target, tvm::IRModule, void, void> const&,
tvm::Target)>::AssignTypedLambda<tvm::$_5>(tvm::$_5,
std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>
>)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}>
>::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs,
tvm::runtime::TVMRetValue*)
118: tvm::TIRToRuntime(tvm::runtime::Map<tvm::Target, tvm::IRModule, void,
void> const&, tvm::Target const&)
117: tvm::codegen::Build(tvm::IRModule, tvm::Target)
116:
tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::runtime::Module
(tvm::IRModule,
tvm::Target)>::AssignTypedLambda<tvm::codegen::$_8>(tvm::codegen::$_8,
std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>
>)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}>
>::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs,
tvm::runtime::TVMRetValue*)
115: tvm::codegen::LLVMModuleNode::Init(tvm::IRModule const&, tvm::Target
const&)
114: void
tvm::codegen::CodeGenLLVM::AddFunctionsOrdered<tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator,
tvm::codegen::CodeGenLLVM::AddFunctionsOrdered<tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator>(tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator, tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void,
void>::iterator)::{lambda(auto:1)#1}>(tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator, tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator,
tvm::codegen::CodeGenLLVM::AddFunctionsOrdered<tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator>(tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator, tvm::runtime::Map<tvm::GlobalVar,
tvm::BaseFunc, void, void>::iterator)::{lambda(auto:1)#1})
113: tvm::codegen::CodeGenCPU::AddFunction(tvm::GlobalVar const&,
tvm::tir::PrimFunc const&)
112: tvm::codegen::CodeGenLLVM::AddFunctionInternal(tvm::GlobalVar const&,
tvm::tir::PrimFunc const&)
111: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
110: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
109: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
108: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
107: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
106: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
105: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
104: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
103: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
102: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AttrStmtNode const*)
101: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AttrStmtNode const*)
100: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
99: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
98: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
97: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
96: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AttrStmtNode const*)
95: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AttrStmtNode const*)
94: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
93: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
92: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
91: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AttrStmtNode const*)
90: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AttrStmtNode const*)
89: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
88: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
87: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
86: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
85: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
84: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
83: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
82: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
81: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
80: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
79: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
78: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
77: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
76: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
75: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
74: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
73: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
72: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
71: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
70: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
69: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
68: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
67: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
66: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
65: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
64: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
63: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
62: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
61: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
60: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
59: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
58: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
57: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
56: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
55: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
54: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
53: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
52: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
51: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
50: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
49: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
48: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
47: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
46: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
45: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
44: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
43: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
42: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
41: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
40: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
39: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
38: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
37: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
36: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
35: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
34: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
33: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
32: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
31: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
30: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
29: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
28: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
27: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
26: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
25: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
24: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
23: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
22: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AssertStmtNode const*)
21: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AssertStmtNode const*)
20: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
19: tvm::codegen::CodeGenCPU::CreateComputeScope(tvm::tir::AttrStmtNode
const*)
18: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
17: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AttrStmtNode const*)
16: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AttrStmtNode const*)
15: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
14: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::LetStmtNode const*)
13: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::AttrStmtNode const*)
12: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AttrStmtNode const*)
11: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
10: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::ForNode const*)
9: tvm::codegen::CodeGenCPU::CreateParallelLaunch(tvm::tir::Stmt const&,
int, std::__cxx11::basic_string<char, std::char_traits<char>,
std::allocator<char> >)
8: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::ForNode const*)
7: tvm::codegen::CodeGenLLVM::CreateSerialFor(llvm::Value*, llvm::Value*,
llvm::Value*, tvm::tir::Var const&, tvm::tir::Stmt const&)
6: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AllocateNode const*)
5: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::AllocateNode const*)
4: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::SeqStmtNode const*)
3: tvm::codegen::CodeGenLLVM::VisitStmt_(tvm::tir::ForNode const*)
2: tvm::codegen::CodeGenLLVM::CreateSerialFor(llvm::Value*, llvm::Value*,
llvm::Value*, tvm::tir::Var const&, tvm::tir::Stmt const&)
1: tvm::codegen::CodeGenCPU::VisitStmt_(tvm::tir::ForNode const*)
0: _ZN3tvm7runtime6detail
File "<path_to_tvm>/tvm/src/target/llvm/codegen_cpu.cc", line 1482
InternalError: Check failed: (is_zero(op->min)) is false:
```
### Triage
Please refer to the list of label tags
[here](https://github.com/apache/tvm/wiki/Issue-Triage-Labels) to find the
relevant tags and add them below in a bullet format (example below).
* needs-triage
* tir
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
