tvm
tvm copied to clipboard
Published
20 hours ago •
apache
apache
[Bug] [TIR] Variable is not defined in the environment IndexError when compiling TIR that looks 'correct'
Steps to Reproduce
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_8(p0: T.Buffer((1, 64, 35, 35), "float32"), p1: T.Buffer((96, 64, 3, 3), "float32"), output_unpack: T.Buffer((1, 96, 35, 35), "float32")):
T.func_attr({"from_legacy_te_schedule": T.bool(True), "hash": "bdd66943e9f8a12c", "target": T.target({"host": {"keys": ["cpu"], "kind": "llvm", "tag": ""}, "keys": ["cpu"], "kind": "llvm", "tag": ""}), "tir.noalias": T.bool(True)})
data_vec = T.allocate([78400], "float32", "global")
data_pad = T.allocate([87616], "float32", "global")
data_vec_1 = T.Buffer((78400,), data=data_vec)
for bs_c_fused_h_fused in T.parallel(560):
for w, vc in T.grid(35, 4):
p0_1 = T.Buffer((78400,), data=p0.data)
data_vec_1[bs_c_fused_h_fused * 140 + w * 4 + vc] = p0_1[bs_c_fused_h_fused // 35 * 4900 + vc * 1225 + bs_c_fused_h_fused % 35 * 35 + w]
data_pad_1 = T.Buffer((87616,), data=data_pad)
for i0_i1_fused_i2_fused in T.parallel(592):
for i3 in range(37):
cse_var_2: T.int32 = i0_i1_fused_i2_fused % 37
cse_var_1: T.int32 = i3 * 4
data_pad_1[i0_i1_fused_i2_fused * 148 + cse_var_1:i0_i1_fused_i2_fused * 148 + cse_var_1 + 4] = T.if_then_else(1 <= cse_var_2 and cse_var_2 < 36 and 1 <= i3 and i3 < 36, data_vec_1[i0_i1_fused_i2_fused // 37 * 4900 + cse_var_2 * 140 + cse_var_1 - 144:i0_i1_fused_i2_fused // 37 * 4900 + cse_var_2 * 140 + cse_var_1 - 144 + 4], T.Broadcast(T.float32(0), 4))
data_vec_2 = T.Buffer((55296,), data=data_vec)
for occ_k_h_fused in T.parallel(72):
for icc, k_w, icb in T.grid(16, 3, 4):
cse_var_4: T.int32 = occ_k_h_fused % 3
cse_var_3: T.int32 = occ_k_h_fused // 3 * 2304
p1_1 = T.Buffer((55296,), 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 + 2304:576]
for n_c_outer_fused_h_fused in T.parallel(840):
conv2d_NCHWc = T.allocate([35], "float32x4", "global")
conv2d_NCHWc_global = T.allocate([7], "float32x4", "global")
conv2d_NCHWc_1 = T.Buffer((35,), "float32x4", data=conv2d_NCHWc)
for ow_outer in range(5):
conv2d_NCHWc_global_1 = T.Buffer((7,), "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)
for ic_outer, kh, kw, ic_inner in T.grid(16, 3, 3, 4):
cse_var_6: T.int32 = n_c_outer_fused_h_fused // 35 * 2304 + ic_outer * 144 + kh * 48 + kw * 16 + ic_inner * 4
cse_var_5: T.int32 = ic_outer * 5476 + kh * 148 + n_c_outer_fused_h_fused % 35 * 148 + ow_outer * 28 + 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]
for ow_inner in range(7):
conv2d_NCHWc_1[ow_outer * 7 + ow_inner] = conv2d_NCHWc_global_1[ow_inner]
for w_outer, w_inner in T.grid(5, 7):
cse_var_7: T.int32 = w_outer * 7 #<- it's defined here?
output_unpack_1 = T.Buffer((117600,), data=output_unpack.data)
b4f = T.float32()
d2e = T.int32()
cb2 = T.int32()
output_unpack_1[T.Let(T.max(T.Broadcast(-2003919016, 4), T.Let(T.Broadcast(1701956818, 4) - (T.Broadcast(-374516067, 4) - T.min(T.Broadcast(-2023512221, 4), T.Broadcast(-795290955, 4)) - T.Broadcast(-408744533, 4)), where={cb2: T.truncmod(cse_var_7, T.Div(cse_var_7, cse_var_7))})) % (T.Cast("int32x4", T.Broadcast(T.uint32(1465149854), 4)) // T.truncmod(T.Broadcast(-1549046547, 4), T.Broadcast(T.Cast("int32", T.acosh(T.Cast("float32", T.Shuffle([T.Cast("int32x4", T.Broadcast(T.uint32(687467356), 4) * T.Broadcast(T.uint32(1937108905), 4))], [3])))), 4))), where={b4f: T.min(T.Let(T.erf(T.min(T.float32(0.80712765069671755) - T.max(T.float32(0.74583677246725477), T.float32(0.65745684962173445)), T.Cast("float32", T.Cast("float32", d2e) + T.float32(0.14202131246354621)) * T.float32(0.3805610012656766))), where={d2e: cse_var_7 % cse_var_7 // (T.max(cse_var_7, cse_var_7) - T.truncmod(cse_var_7, cse_var_7))}), T.float32(0.75022262892839819))})] = conv2d_NCHWc_1[cse_var_7 + w_inner]
func = tvmgen_default_fused_nn_conv2d_8
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)
Expected Behavior
At the very least, cse_var_7 should be defined, so I shouldn't get error messages about it not being defined.
Observed Behavior
(Note, I've updated file paths in this log to simplify paths to TVM and where I ran the script above)
[21:04:21] <path_to_tvm>tvm/src/script/printer/tir/expr.cc:70: Warning: Didn't find variable definition for: cse_var_7
[21:04:21] <path_to_tvm>tvm/src/script/printer/ir/../utils.h:46: Warning: TVMScript printer falls back to the legacy ReprPrinter with the error:
[21:04:21] <path_to_tvm>tvm/src/script/printer/tir/expr.cc:76: IndexError: Variable is not defined in the environment: cse_var_7
Stack trace:
0: _ZN3tvm7runtime6detail
1: tvm::script::printer::PrintVar(tvm::tir::Var const&, tvm::ObjectPath const&, tvm::script::printer::IRDocsifier const&)
2: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::script::printer::Doc (tvm::tir::Var, tvm::ObjectPath, tvm::script::printer::IRDocsifier)>::AssignTypedLambda<tvm::script::printer::$_0>(tvm::script::printer::$_0)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
3: tvm::runtime::TVMRetValue tvm::runtime::PackedFunc::operator()<tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&>(tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&) const
4: tvm::script::printer::Doc tvm::script::printer::IRDocsifierFunctor<tvm::script::printer::Doc, tvm::ObjectPath, tvm::script::printer::IRDocsifier>::operator()<tvm::runtime::ObjectRef>(tvm::runtime::String const&, tvm::runtime::ObjectRef, tvm::ObjectPath, tvm::script::printer::IRDocsifier) const
5: tvm::script::printer::ExprDoc tvm::script::printer::IRDocsifierNode::AsDoc<tvm::script::printer::ExprDoc>(tvm::runtime::ObjectRef const&, tvm::ObjectPath const&) const
6: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::script::printer::Doc (tvm::tir::FloorMod, tvm::ObjectPath, tvm::script::printer::IRDocsifier)>::AssignTypedLambda<tvm::script::printer::$_21>(tvm::script::printer::$_21)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
7: tvm::runtime::TVMRetValue tvm::runtime::PackedFunc::operator()<tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&>(tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&) const
8: tvm::script::printer::Doc tvm::script::printer::IRDocsifierFunctor<tvm::script::printer::Doc, tvm::ObjectPath, tvm::script::printer::IRDocsifier>::operator()<tvm::runtime::ObjectRef>(tvm::runtime::String const&, tvm::runtime::ObjectRef, tvm::ObjectPath, tvm::script::printer::IRDocsifier) const
9: tvm::script::printer::ExprDoc tvm::script::printer::IRDocsifierNode::AsDoc<tvm::script::printer::ExprDoc>(tvm::runtime::ObjectRef const&, tvm::ObjectPath const&) const
10: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::script::printer::Doc (tvm::tir::FloorDiv, tvm::ObjectPath, tvm::script::printer::IRDocsifier)>::AssignTypedLambda<tvm::script::printer::$_20>(tvm::script::printer::$_20)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
11: tvm::runtime::TVMRetValue tvm::runtime::PackedFunc::operator()<tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&>(tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&) const
12: tvm::script::printer::Doc tvm::script::printer::IRDocsifierFunctor<tvm::script::printer::Doc, tvm::ObjectPath, tvm::script::printer::IRDocsifier>::operator()<tvm::runtime::ObjectRef>(tvm::runtime::String const&, tvm::runtime::ObjectRef, tvm::ObjectPath, tvm::script::printer::IRDocsifier) const
13: tvm::script::printer::ExprDoc tvm::script::printer::IRDocsifierNode::AsDoc<tvm::script::printer::ExprDoc>(tvm::runtime::ObjectRef const&, tvm::ObjectPath const&) const
14: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::script::printer::Doc (tvm::Range, tvm::ObjectPath, tvm::script::printer::IRDocsifier)>::AssignTypedLambda<tvm::script::printer::$_13>(tvm::script::printer::$_13)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
15: tvm::runtime::TVMRetValue tvm::runtime::PackedFunc::operator()<tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&>(tvm::runtime::ObjectRef&, tvm::ObjectPath&, tvm::script::printer::IRDocsifier&) const
16: tvm::script::printer::Doc tvm::script::printer::IRDocsifierFunctor<tvm::script::printer::Doc, tvm::ObjectPath, tvm::script::printer::IRDocsifier>::operator()<tvm::runtime::ObjectRef>(tvm::runtime::String const&, tvm::runtime::ObjectRef, tvm::ObjectPath, tvm::script::printer::IRDocsifier) const
17: tvm::script::printer::Doc tvm::script::printer::IRDocsifierNode::AsDoc<tvm::script::printer::Doc>(tvm::runtime::ObjectRef const&, tvm::ObjectPath const&) const
18: tvm::script::printer::Docsify[abi:cxx11](tvm::runtime::ObjectRef const&, tvm::script::printer::IRDocsifier const&, tvm::script::printer::Frame const&, tvm::PrinterConfig const&)
19: tvm::script::printer::ReprPrintIR[abi:cxx11](tvm::runtime::ObjectRef const&, tvm::PrinterConfig const&)
20: tvm::NodeFunctor<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > (tvm::runtime::ObjectRef const&, tvm::PrinterConfig const&)>::operator()(tvm::runtime::ObjectRef const&, tvm::PrinterConfig const&) const
21: tvm::TVMScriptPrinter::Script[abi:cxx11](tvm::runtime::ObjectRef const&, tvm::runtime::Optional<tvm::PrinterConfig> const&)
22: tvm::script::printer::RedirectedReprPrinterMethod(tvm::runtime::ObjectRef const&, tvm::ReprPrinter*)
23: tvm::runtime::operator<<(std::ostream&, tvm::runtime::ObjectRef const&)
24: tvm::arith::TransitiveComparisonAnalyzer::Impl::Bind(tvm::tir::Var const&, tvm::Range const&, bool)
25: tvm::arith::TransitiveComparisonAnalyzer::Impl::Bind(tvm::tir::Var const&, tvm::PrimExpr const&, bool)
26: tvm::arith::Analyzer::Bind(tvm::tir::Var const&, tvm::PrimExpr const&, bool)
27: tvm::arith::RewriteSimplifier::Impl::VisitExpr_(tvm::tir::LetNode const*)
28: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl10VisitExp
29: _ZZN3tvm3tir11ExprFunctorIFNS_8PrimExprERKS2_EE10I
30: tvm::NodeFunctor<tvm::PrimExpr (tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*)>::operator()(tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*) const
31: tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>::VisitExpr(tvm::PrimExpr const&)
32: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl9VisitExpr
33: tvm::arith::RewriteSimplifier::Impl::VisitExpr_(tvm::tir::MinNode const*)
34: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl10VisitExp
35: _ZZN3tvm3tir11ExprFunctorIFNS_8PrimExprERKS2_EE10Init
36: tvm::NodeFunctor<tvm::PrimExpr (tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*)>::operator()(tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*) const
37: tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>::VisitExpr(tvm::PrimExpr const&)
38: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl9VisitExpr
39: tvm::arith::RewriteSimplifier::operator()(tvm::PrimExpr const&)
40: tvm::arith::Analyzer::Bind(tvm::tir::Var const&, tvm::PrimExpr const&, bool)
41: non-virtual thunk to tvm::arith::IRVisitorWithAnalyzer::VisitExpr_(tvm::tir::LetNode const*)
42: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::BufferStoreNode const*)
43: tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tir::ForNode const*)
44: tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tir::ForNode const*)
45: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
46: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
47: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
48: tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tir::ForNode const*)
49: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
50: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
51: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
52: tvm::tir::TextureFlatten(tvm::tir::PrimFunc)
53: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::tir::PrimFunc (tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)>::AssignTypedLambda<tvm::tir::transform::TextureFlatten()::$_0>(tvm::tir::transform::TextureFlatten()::$_0)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
54: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
55: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
56: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
57: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
58: tvm::transform::Pass::operator()(tvm::IRModule) const
59: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>)
60: tvm::LowerModule(tvm::IRModule, bool)
61: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::IRModule (tvm::IRModule, bool)>::AssignTypedLambda<tvm::$_2>(tvm::$_2, 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*)
Traceback (most recent call last):
File "<path_to_bug_reproduction>/TVMBugReport1/reprod.py", line 70, in <module>
nopt_mod = tvm.build(mod)
File "<path_to_tvm>tvm/python/tvm/driver/build_module.py", line 239, in build
input_mod = lower(inputs)
File "<path_to_tvm>tvm/python/tvm/driver/build_module.py", line 130, in lower
return ffi.lower_module(inp, simple_mode)
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
IndexError: Traceback (most recent call last):
38: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::IRModule (tvm::IRModule, bool)>::AssignTypedLambda<tvm::$_2>(tvm::$_2, 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*)
37: tvm::LowerModule(tvm::IRModule, bool)
36: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>)
35: tvm::transform::Pass::operator()(tvm::IRModule) const
34: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
33: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
32: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
31: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
30: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::tir::PrimFunc (tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)>::AssignTypedLambda<tvm::tir::transform::TextureFlatten()::$_0>(tvm::tir::transform::TextureFlatten()::$_0)::{lambda(tvm::runtime::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
29: tvm::tir::TextureFlatten(tvm::tir::PrimFunc)
28: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
27: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
26: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
25: tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tir::ForNode const*)
24: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
23: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::AllocateNode const*)
22: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
21: tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tir::ForNode const*)
20: tvm::arith::IRVisitorWithAnalyzer::VisitStmt_(tvm::tir::ForNode const*)
19: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::BufferStoreNode const*)
18: non-virtual thunk to tvm::arith::IRVisitorWithAnalyzer::VisitExpr_(tvm::tir::LetNode const*)
17: tvm::arith::Analyzer::Bind(tvm::tir::Var const&, tvm::PrimExpr const&, bool)
16: tvm::arith::RewriteSimplifier::operator()(tvm::PrimExpr const&)
15: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl9VisitExpr
14: tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>::VisitExpr(tvm::PrimExpr const&)
13: tvm::NodeFunctor<tvm::PrimExpr (tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*)>::operator()(tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*) const
12: _ZZN3tvm3tir11ExprFunctorIFNS_8PrimExprERKS2_EE10Init
11: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl10VisitExp
10: tvm::arith::RewriteSimplifier::Impl::VisitExpr_(tvm::tir::MinNode const*)
9: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl9VisitExpr
8: tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>::VisitExpr(tvm::PrimExpr const&)
7: tvm::NodeFunctor<tvm::PrimExpr (tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*)>::operator()(tvm::runtime::ObjectRef const&, tvm::tir::ExprFunctor<tvm::PrimExpr (tvm::PrimExpr const&)>*) const
6: _ZZN3tvm3tir11ExprFunctorIFNS_8PrimExprERKS2_EE10I
5: _ZThn16_N3tvm5arith17RewriteSimplifier4Impl10VisitExp
4: tvm::arith::RewriteSimplifier::Impl::VisitExpr_(tvm::tir::LetNode const*)
3: tvm::arith::Analyzer::Bind(tvm::tir::Var const&, tvm::PrimExpr const&, bool)
2: tvm::arith::TransitiveComparisonAnalyzer::Impl::Bind(tvm::tir::Var const&, tvm::PrimExpr const&, bool)
1: tvm::arith::TransitiveComparisonAnalyzer::Impl::Bind(tvm::tir::Var const&, tvm::Range const&, bool)
0: _ZN3tvm7runtime6detail
File "<path_to_tvm>tvm/src/script/printer/tir/expr.cc", line 76
IndexError: Variable is not defined in the environment: cse_var_7range(min=floordiv(floormod(cse_var_7, cse_var_7), (cse_var_7 - (cse_var_7 % cse_var_7))), ext=1)Range(0x55c18f0b6240)
Additoinal notes
Printing the IR before and after each pass seems to indicate that the tir.TextureFlatten pass is the source of the exception. The before pass IR is printed, but no after pass IR is printed.
Triage
- needs-triage
- tir
cc @Hzfengsy @junrushao @quic-sanirudh @shingjan
