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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
