Error during build process with custom operators: Bind have an unmet assertion

Necrotos · June 3, 2023, 6:23pm

Hello, I am trying to generate C code for a fairly complex model in order to run on my complex accelerator. During the build process, I get the following error:

Traceback (most recent call last):
   File "/home/git/tvm/python/tvm/relay/build_module.py", line 364, in build
        graph_json, runtime_mod, params = bld_mod.build(
   File "/home/git/tvm/python/tvm/relay/build_module.py", line 161, in build
        self._build(
   File "/home/git/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line 237, in __call__
        raise get_last_ffi_error()
tvm._ffi.base.TVMError: Traceback (most recent call last):
  88: TVMFuncCall
  87: tvm::relay::backend::RelayBuildModule::GetFunction(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#3}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const
  86: tvm::relay::backend::RelayBuildModule::BuildRelay(tvm::IRModule, tvm::runtime::String const&)
  85: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::relay::backend::AOTExecutorCodegenModule::GetFunction(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
  84: tvm::relay::backend::AOTExecutorCodegen::Codegen(tvm::IRModule, tvm::relay::Function, tvm::runtime::String)
  83: tvm::transform::Pass::operator()(tvm::IRModule) const
  82: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  81: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  80: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  79: tvm::transform::ModulePassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  78: _ZN3tvm7runtime13PackedFuncObj9ExtractorINS0_16PackedFuncSubObjIZNS0_15TypedPackedFuncIFNS_8IRModuleES5_NS_9transform11PassContextEEE17AssignTypedLambdaIZNS_5relay3tec7LowerTEENS0_6StringENS_17CompilationConfigESt8functionIFvNS_8BaseFuncEEEEUlS5_S7_E_EEvT_EUlRKNS0_7TVMArgsEPNS0_11TVMRetValueEE_EEE4CallEPKS1_SL_SP_
  77: tvm::relay::tec::LowerTE(tvm::IRModule const&, tvm::runtime::String const&, std::function<void (tvm::BaseFunc)>, tvm::CompilationConfig)
  76: tvm::transform::Pass::operator()(tvm::IRModule) const
  75: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  74: tvm::relay::transform::FunctionPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  73: _ZN3tvm7runtime13PackedFuncObj9ExtractorINS0_16PackedFuncSubObjIZNS0_15TypedPackedFuncIFNS_5relay8FunctionES6_NS_8IRModuleENS_9transform11PassContextEEE17AssignTypedLambdaIZNS5_3tec15LowerTensorExprENSD_10TECompilerESt8functionIFvNS_8BaseFuncEEENS_17CompilationConfigEEUlS6_S7_S9_E_EEvT_EUlRKNS0_7TVMArgsEPNS0_11TVMRetValueEE_EEE4CallEPKS1_SM_SQ_
  72: tvm::relay::ExprMutator::VisitExpr(tvm::RelayExpr const&)
  71: tvm::relay::ExprFunctor<tvm::RelayExpr (tvm::RelayExpr const&)>::VisitExpr(tvm::RelayExpr const&)
  70: _ZZN3tvm5relay11ExprFunc
  69: tvm::relay::transform::DeviceAwareExprMutator::VisitExpr_(tvm::relay::FunctionNode const*)
  68: tvm::relay::tec::LowerTensorExprMutator::DeviceAwareVisitExpr_(tvm::relay::FunctionNode const*)
  67: _ZN3tvm5relay9
  66: tvm::relay::ExprMutator::VisitExpr_(tvm::relay::FunctionNode const*)
  65: tvm::relay::ExprMutator::VisitExpr(tvm::RelayExpr const&)
  64: tvm::relay::ExprFunctor<tvm::RelayExpr (tvm::RelayExpr const&)>::VisitExpr(tvm::RelayExpr const&)
  63: _ZZN3tvm5relay11ExprFunc
  62: tvm::relay::transform::DeviceAwareExprMutator::VisitExpr_(tvm::relay::LetNode const*)
  61: tvm::relay::tec::LowerTensorExprMutator::PreVisitLetBinding_(tvm::relay::Var const&, tvm::RelayExpr const&)
  60: tvm::relay::ExprMutator::VisitExpr(tvm::RelayExpr const&)
  59: tvm::relay::ExprFunctor<tvm::RelayExpr (tvm::RelayExpr const&)>::VisitExpr(tvm::RelayExpr const&)
  58: _ZZN3tvm5relay11ExprFunc
  57: tvm::relay::transform::DeviceAwareExprMutator::VisitExpr_(tvm::relay::CallNode const*)
  56: tvm::relay::ExprMutator::VisitExpr(tvm::RelayExpr const&)
  55: tvm::relay::ExprFunctor<tvm::RelayExpr (tvm::RelayExpr const&)>::VisitExpr(tvm::RelayExpr const&)
  54: _ZZN3tvm5relay11ExprFunc
  53: tvm::relay::transform::DeviceAwareExprMutator::VisitExpr_(tvm::relay::CallNode const*)
  52: tvm::relay::tec::LowerTensorExprMutator::DeviceAwareVisitExpr_(tvm::relay::CallNode const*)
  51: tvm::relay::tec::TECompilerImpl::Lower(tvm::relay::tec::CCacheKey const&)
  50: tvm::relay::tec::TECompilerImpl::LowerInternal(tvm::relay::tec::CCacheKey const&, tvm::GlobalVarSupply)
  49: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::te::Tensor, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, tvm::GlobalVarSupply, bool)
  48: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, tvm::GlobalVarSupply, bool)
  47: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>)
  46: tvm::transform::Pass::operator()(tvm::IRModule) const
  45: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  44: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  43: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  42: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  41: _ZN3tvm7runtime13PackedFuncObj9ExtractorINS0_16PackedFuncSubObjIZNS0_15TypedPackedFuncIFNS_3tir8PrimFuncES6_NS_8IRModuleENS_9transform11PassContextEEE17AssignTypedLambdaIZNS5_9transform14StorageFlattenEibEUlS6_S7_S9_E_EEvT_EUlRKNS0_7TVMArgsEPNS0_11TVMRetValueEE_EEE4CallEPKS1_SG_SK_
  40: tvm::tir::StorageFlatten(tvm::tir::PrimFunc, int, bool)
  39: tvm::transform::Pass::operator()(tvm::IRModule) const
  38: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  37: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  36: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  35: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
  34: _ZN3tvm7runtime13PackedFuncObj9ExtractorINS0_16PackedFuncSubObjIZNS0_15TypedPackedFuncIFNS_3tir8PrimFuncES6_NS_8IRModuleENS_9transform11PassContextEEE17AssignTypedLambdaIZNS5_19BufferShapeLegalize4PassEvEUlS6_S7_S9_E_EEvT_EUlRKNS0_7TVMArgsEPNS0_11TVMRetValueEE_EEE4CallEPKS1_SG_SK_
  33: tvm::tir::BufferShapeLegalize::Pass()::{lambda(tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)#1}::operator()(tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext) const
  32: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  31: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  30: _ZZN3tvm3tir11StmtFuncto
  29: tvm::tir::BufferShapeLegalize::VisitStmt_(tvm::tir::BufferRealizeNode const*)
  28: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  27: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  26: _ZZN3tvm3tir11StmtFunctorIF
  25: tvm::tir::StmtMutator::VisitStmt_(tvm::tir::SeqStmtNode const*)
  24: tvm::runtime::ObjectPtr<tvm::runtime::Object> tvm::runtime::Array<tvm::tir::Stmt, void>::MapHelper<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::runtime::Object, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
  23: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  22: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  21: _ZZN3tvm3tir11StmtFuncto
  20: tvm::tir::BufferShapeLegalize::VisitStmt_(tvm::tir::BufferRealizeNode const*)
  19: tvm::tir::StmtMutator::VisitStmt_(tvm::tir::BufferRealizeNode const*)
  18: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  17: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  16: _ZZN3tvm3tir11StmtFuncto
  15: tvm::tir::BufferShapeLegalize::VisitStmt_(tvm::tir::AttrStmtNode const*)
  14: tvm::tir::StmtMutator::VisitStmt_(tvm::tir::AttrStmtNode const*)
  13: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  12: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  11: _ZZN3tvm3tir11StmtFuncto
  10: tvm::tir::BufferShapeLegalize::VisitStmt_(tvm::tir::AttrStmtNode const*)
  9: tvm::tir::StmtMutator::VisitStmt_(tvm::tir::AttrStmtNode const*)
  8: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
  7: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
  6: _ZZN3tvm3tir11StmtFuncto
  5: tvm::tir::BufferShapeLegalize::VisitStmt_(tvm::tir::AttrStmtNode const*)
  4: tvm::tir::BufferShapeLegalize::HandleBufferBindScope(tvm::tir::AttrStmtNode const*)
  3: tvm::tir::ArgBinder::BindBuffer(tvm::tir::Buffer const&, tvm::tir::Buffer const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool)
  2: tvm::tir::ArgBinder::BindArray(tvm::runtime::Array<tvm::PrimExpr, void> const&, tvm::runtime::Array<tvm::PrimExpr, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&)
  1: tvm::tir::ArgBinder::Bind_(tvm::PrimExpr const&, tvm::PrimExpr const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool)
  0: tvm::tir::BinderAddAssert(tvm::arith::Analyzer*, tvm::PrimExpr, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<tvm::tir::Stmt, std::allocator<tvm::tir::Stmt> >*)
tvm._ffi.base.TVMError: Traceback (most recent call last):
TVMError: Bind have an unmet assertion: T.bool(False),  on argument inp_buff.shape[1] scond: T.bool(False)

Is there some way to get more information on what exactly is causing this issue? I don’t understand where exactly arg_binder comes into play and the error isn’t providing that much information for me.

Also, I’d be interested in how to read this stack trace. I don’t understand what the numbers represent. Thank you!

Necrotos · June 6, 2023, 12:23pm

I spent some time on this issue. It is related to the size of the input buffer of my conv2d. Where are the shapes passed to ArgBinder? I can’t find where exactly this happens and where the values come from.

It might be related to the way decl_buffer is called for the input. I do it like this:

inp_layout = tvm.tir.decl_buffer(
        inp_shape,
        env.inp_dtype,
        "inp_buff",
        strides=[te.var("inp_x"), te.var("inp_y"), te.var("inp_b"), te.var("inp_k")],
    )

Necrotos · June 6, 2023, 10:10pm

I have added a LOG info into arg_binder. The error shows up in ArgBinder::BindBuffer when value to arg. With the additional info I have gathered I can see this:

[00:03:00] /home/git/tvm/src/tir/transforms/arg_binder.cc:146: ArgBinder::BindBuffer #############----arg_name: inp_buff
[00:03:00] /home/git/tvm/src/tir/transforms/arg_binder.cc:147: ArgBinder::BindBuffer #############----value_name: inp_buff
[00:03:00] /home/git/tvm/src/tir/transforms/arg_binder.cc:148: ArgBinder::BindBuffer #############----arg: inp_buff,  value inp_buff, arg_name inp_buff arg->shape [1, 514, 258, 128] value->shape [1, 513, 257, 128]
[00:03:00] /home/git/tvm/src/tir/transforms/arg_binder.cc:83: ArgBinder::BindArray #############----arg: [1, 514, 258, 128],  value [1, 513, 257, 128], arg_name inp_buff.shape

Clearly, value has an incorrect shape here. Where does this value come from? I am unable to find this information, as I can’t find anything that matches this shape anywhere in my code. I have searched in the declaration of the buffers, the definition of the intrin function, the original PyTorch model, but I can’t find this specific shape anywhere.