Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[Bug] [TIR] Variable is not defined in the environment IndexError when compiling TIR that looks 'correct' #17387

Open
talha-ahsan opened this issue Sep 19, 2024 · 0 comments
Labels
needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it type: bug

Comments

@talha-ahsan
Copy link

talha-ahsan commented Sep 19, 2024

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

@talha-ahsan talha-ahsan added needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it type: bug labels Sep 19, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it type: bug
Projects
None yet
Development

No branches or pull requests

1 participant