apache / apache/tvm

[Bug] [CodeGen] InternalError: Check failed: (is_zero(op->min)) is false during CodeGen

Open
#17,386 1 comment 0 reactions 0 assignees View on GitHub
needs-triage type: bug
Dominant language
Python
Stars
13.7k
Forks
4k
Avg merge
2d 19h
Merged PRs (30d)
111

Description

### 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 "/TVMBugReport3/reprod.py", line 113, in
nopt_mod = tvm.build(mod)
File "/tvm/python/tvm/driver/build_module.py", line 297, in build
rt_mod_host = _driver_ffi.tir_to_runtime(annotated_mods, target_host)
File "/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line 239, in __call__
raise_last_ffi_error()
File "/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 const&, tvm::Target)>::AssignTypedLambda(tvm::$_5, std::__cxx11::basic_string, std::allocator >)::{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 const&, tvm::Target const&)
117: tvm::codegen::Build(tvm::IRModule, tvm::Target)
116: tvm::runtime::PackedFuncObj::Extractor::AssignTypedLambda(tvm::codegen::$_8, std::__cxx11::basic_string, std::allocator >)::{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::iterator, tvm::codegen::CodeGenLLVM::AddFunctionsOrdered::iterator>(tvm::runtime::Map::iterator, tvm::runtime::Map::iterator)::{lambda(auto:1)#1}>(tvm::runtime::Map::iterator, tvm::runtime::Map::iterator, tvm::codegen::CodeGenLLVM::AddFunctionsOrdered::iterator>(tvm::runtime::Map::iterator, tvm::runtime::Map::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, std::allocator >)
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 "/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

Contributor guide

No contributing guide indexed for this repository

Research direction

Start with the reproducer in reprod.py and run the tvm.build(mod) call under the shown PassContext. Trace the failure from tir.LowerIntrin through TIRToRuntime and codegen Build, using the supplied TIR and InternalError as the test case. Done means the reproducer no longer reaches the opaque check failure and instead compiles or reports which invalid TIR condition caused it.

Written by the indexing model from the issue text.

Assessment

Tech stack
python
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.