You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "Cookiee235 (via GitHub)" <gi...@apache.org> on 2025/05/09 13:27:45 UTC
[I] [Bug] InternalError: Check failed: (ret_ex.dtype() == var.dtype()) is false [tvm]
Cookiee235 opened a new issue, #17937:
URL: https://github.com/apache/tvm/issues/17937
### Actual behavior
```
Traceback (most recent call last):
File "/data/qshenaf/remote_pc/TirFuzz/bugs/05-03_20-50/topi.nn.group_norm_4.py", line 20, in <module>
tvm.build(opt_mod, target='llvm')
File "/data/qshenaf/envs/tvm/python/tvm/driver/build_module.py", line 59, in build
return tvm.tir.build(mod, target, pipeline)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/python/tvm/tir/build.py", line 173, in build
mod = pipeline(mod)
^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/python/tvm/ir/transform.py", line 238, in __call__
return _ffi_transform_api.RunPass(self, mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "tvm/_ffi/_cython/./packed_func.pxi", line 339, in tvm._ffi._cy3.core.PackedFuncBase.__call__
File "tvm/_ffi/_cython/./packed_func.pxi", line 270, in tvm._ffi._cy3.core.FuncCall
File "tvm/_ffi/_cython/./packed_func.pxi", line 259, in tvm._ffi._cy3.core.FuncCall3
File "tvm/_ffi/_cython/./base.pxi", line 185, in tvm._ffi._cy3.core.CHECK_CALL
File "/data/qshenaf/envs/tvm/python/tvm/_ffi/base.py", line 468, in raise_last_ffi_error
raise py_err
File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback
File "/data/qshenaf/envs/tvm/python/tvm/tir/pipeline.py", line 122, in _pipeline
mod = tvm.ir.transform.Sequential(passes)(mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/python/tvm/ir/transform.py", line 238, in __call__
return _ffi_transform_api.RunPass(self, mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "tvm/_ffi/_cython/./packed_func.pxi", line 339, in tvm._ffi._cy3.core.PackedFuncBase.__call__
File "tvm/_ffi/_cython/./packed_func.pxi", line 270, in tvm._ffi._cy3.core.FuncCall
File "tvm/_ffi/_cython/./packed_func.pxi", line 259, in tvm._ffi._cy3.core.FuncCall3
File "tvm/_ffi/_cython/./base.pxi", line 185, in tvm._ffi._cy3.core.CHECK_CALL
File "/data/qshenaf/envs/tvm/src/tir/ir/transform.cc", line 121, in tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
func = pass_func(std::move(func), mod, pass_ctx);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 255, in operator()
return PlanAndUpdateBufferAllocationLocation(std::move(f));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 247, in tvm::tir::PlanAndUpdateBufferAllocationLocation(tvm::tir::PrimFunc)
fptr->body = locator(fptr->body);
^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 178, in tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::BlockNode const*)
Stmt stmt = StmtMutator::VisitStmt_(op);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 211, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)
return MutateArray(self, arr, fmutate);
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 184, in tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void>const&)::{lambda(tvm::tir::Stmt const&)#1})
Array<T> copy = arr.Map(fmutate);
^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 652, in tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::runtime::Array<tvm::tir::Stmt, void>::Map<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::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}) const
return Array<U>(MapHelper(data_, fmap));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 877, in 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::ObjectPtr<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})
U mapped = fmap(DowncastNoCheck<T>(*it));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 210, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
auto fmutate = [self](const Stmt& s) { return self->VisitStmt(s); };
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 145, in tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
auto node = Downcast<For>(StmtMutator::VisitStmt_(op));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 211, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)
return MutateArray(self, arr, fmutate);
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 184, in tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void>const&)::{lambda(tvm::tir::Stmt const&)#1})
Array<T> copy = arr.Map(fmutate);
^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 652, in tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::runtime::Array<tvm::tir::Stmt, void>::Map<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::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}) const
return Array<U>(MapHelper(data_, fmap));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 877, in 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::ObjectPtr<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})
U mapped = fmap(DowncastNoCheck<T>(*it));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 210, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
auto fmutate = [self](const Stmt& s) { return self->VisitStmt(s); };
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 156, in tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
node.CopyOnWrite()->body = InjectOpaqueBlock(node->body, new_block_alloc_bufs);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 219, in tvm::tir::BufferAllocationLocator::InjectOpaqueBlock(tvm::tir::Stmt, tvm::runtime::Array<tvm::tir::Buffer, void> const&)
GetBlockReadWriteRegion(opaque_block, buffer_data_to_buffer_);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 385, in tvm::tir::GetBlockReadWriteRegion(tvm::tir::Block const&, tvm::runtime::Map<tvm::tir::Var, tvm::tir::Buffer, void, void> const&)
detector(block);
^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 135, in tvm::tir::BlockReadWriteDetector::operator()(tvm::tir::Stmt const&)
StmtExprVisitor::operator()(stmt);
^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 144, in tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::BlockNode const*)
this->VisitStmt(op->body);
^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 119, in tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
VisitArray(op->seq, [this](const Stmt& s) { this->VisitStmt(s); });
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/functor_common.h", line 35, in VisitArray<tvm::tir::Stmt, tvm::tir::StmtVisitor::VisitStmt_(const tvm::tir::SeqStmtNode*)::<lambda(const tvm::tir::Stmt&)> >
fvisit(arr[i]);
^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 119, in operator()
VisitArray(op->seq, [this](const Stmt& s) { this->VisitStmt(s); });
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 167, in tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::ForNode const*)
StmtVisitor::VisitStmt_(op);
^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 48, in tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::ForNode const*)
this->VisitStmt(op->body);
^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 258, in tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::BlockRealizeNode const*)
Substitute(range->min, vmap), Substitute(range->extent, vmap))),
^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 758, in tvm::tir::Substitute(tvm::PrimExpr, std::function<tvm::runtime::Optional<tvm::PrimExpr> (tvm::tir::Var const&)>)
return IRSubstitute(vmap)(std::move(expr));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 660, in tvm::tir::IRSubstitute::VisitExpr_(tvm::tir::VarNode const*)
ICHECK(ret_ex.dtype() == var.dtype()) << "substituting " << var << ":" << var.dtype()
^^^^^^^^^^^^^^^^^^^^^^^^^^^
tvm.error.InternalError: Traceback (most recent call last):
37: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/transform.cc:121
36: operator()
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:255
35: tvm::tir::PlanAndUpdateBufferAllocationLocation(tvm::tir::PrimFunc)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:247
34: tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::BlockNode const*)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:178
33: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:211
32: tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:184
31: tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::runtime::Array<tvm::tir::Stmt, void>::Map<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::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}) const
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:652
30: 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::ObjectPtr<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})
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:877
29: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:210
28: tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:145
27: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:211
26: tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1})
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:184
25: tvm::runtime::Array<tvm::tir::Stmt, std::enable_if<std::is_base_of<tvm::runtime::ObjectRef, tvm::tir::Stmt>::value, void>::type> tvm::runtime::Array<tvm::tir::Stmt, void>::Map<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::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}) const
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:652
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::ObjectPtr<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})
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:877
23: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:210
22: tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:156
21: tvm::tir::BufferAllocationLocator::InjectOpaqueBlock(tvm::tir::Stmt, tvm::runtime::Array<tvm::tir::Buffer, void> const&)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:219
20: tvm::tir::GetBlockReadWriteRegion(tvm::tir::Block const&, tvm::runtime::Map<tvm::tir::Var, tvm::tir::Buffer, void, void> const&)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:385
19: tvm::tir::BlockReadWriteDetector::operator()(tvm::tir::Stmt const&)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:135
18: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::BlockNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:144
17: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:119
16: VisitArray<tvm::tir::Stmt, tvm::tir::StmtVisitor::VisitStmt_(const tvm::tir::SeqStmtNode*)::<lambda(const tvm::tir::Stmt&)> >
at /data/qshenaf/envs/tvm/src/tir/ir/functor_common.h:35
15: operator()
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:119
14: tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:167
13: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:48
12: tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::BlockRealizeNode const*)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:258
11: tvm::tir::Substitute(tvm::PrimExpr, std::function<tvm::runtime::Optional<tvm::PrimExpr> (tvm::tir::Var const&)>)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:758
10: _ZThn16_N3tvm3tir15StmtExprMu
9: _ZThn16_N3tvm3tir15StmtExprMu
8: _ZThn16_N3tvm3tir15StmtExprMu
7: _ZThn16_N3tvm3tir15StmtExprMu
6: _ZThn16_N3tvm3tir15StmtExprMu
5: _ZThn16_N3tvm3tir15StmtExprMu
4: _ZThn16_N3tvm3tir15StmtExprMu
3: _ZThn16_N3tvm3tir15StmtExprMu
2: _ZThn16_N3tvm3tir15StmtExprMu
1: _ZThn16_N3tvm3tir12IRSubstitu
0: tvm::tir::IRSubstitute::VisitExpr_(tvm::tir::VarNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:660
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 660
InternalError: Check failed: (ret_ex.dtype() == var.dtype()) is false: substituting v_ax4:int32 -> T.Ramp(0, 1, 3):int32x3
```
### Environment
tvm-0.21.dev0
### Steps to reproduce
```
import tvm
from tvm import te, topi, tir
from tvm import meta_schedule as ms
data = te.placeholder((2, 48, 28, 28, 3), dtype='float16', name='data')
gamma = te.placeholder((48,), dtype='float16', name='gamma')
beta = te.placeholder((48,), dtype='float16', name='beta')
op_config = {'data': data, 'gamma': gamma, 'beta': beta, 'num_groups': 12, 'channel_axis': 1, 'axes': [2, 3, 4], 'epsilon': 1e-08, }
op_output = topi.nn.group_norm(**op_config)
sch = tir.Schedule(te.create_prim_func([data, gamma, beta, op_output]).with_attr('target', tvm.target.Target('llvm')))
database = ms.tir_integration.tune_tir(sch.mod, target='llvm --num-cores=16', work_dir='./tune_tmp', max_trials_global=1, num_trials_per_iter=1)
sch = ms.tir_integration.compile_tir(database, sch.mod, 'llvm --num-cores=16')
passes = [tir.transform.VectorizeLoop(True),tir.transform.InjectVirtualThread()]
with tvm.ir.transform.PassContext(opt_level=4):
opt_mod = tvm.ir.transform.Sequential(passes)(sch.mod)
tvm.build(opt_mod, target='llvm')
```
### Triage
* needs-triage
* tir
* tune:meta_schedule
--
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: commits-unsubscribe@tvm.apache.org.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org