Cookiee235 opened a new issue, #17936:
URL: https://github.com/apache/tvm/issues/17936
### 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 19,
tvm.build(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/compact_buffer_region.cc",
line 755, in operator()
return CompactBufferAllocation(std::move(f), is_strict);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 745, in tvm::tir::CompactBufferAllocation(tvm::tir::PrimFunc, bool)
auto region = BufferAccessRegionCollector::Collect(f,
/*collect_inbound=*/is_strict);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 114, in tvm::tir::BufferAccessRegionCollector::Collect(tvm::tir::PrimFunc
const&, bool)
region_collector(f->body);
^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 314, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::AllocateNode const*)
StmtExprVisitor::VisitStmt(op->body);
^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 314, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::AllocateNode const*)
StmtExprVisitor::VisitStmt(op->body);
^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 314, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::AllocateNode const*)
StmtExprVisitor::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/transforms/compact_buffer_region.cc",
line 159, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::ForNode const*)
StmtExprVisitor::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/transforms/compact_buffer_region.cc",
line 170, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::LetStmtNode const*)
StmtExprVisitor::VisitStmt(op->body);
^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 170, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::LetStmtNode const*)
StmtExprVisitor::VisitStmt(op->body);
^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 134, in
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::BufferStoreNode
const*)
VisitBufferAccess(BufferRegion::FromPoint(op->buffer, op->indices));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 377, in
tvm::tir::BufferAccessRegionCollector::VisitBufferAccess(tvm::tir::BufferRegion
const&)
NDIntSetEval(buffer_region->region, predicate, dom_map_, &dom_analyzer_);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc",
line 54, in tvm::tir::NDIntSetEval(tvm::runtime::Array<tvm::Range, void>,
tvm::PrimExpr, std::unordered_map<tvm::tir::VarNode const*, tvm::arith::IntSet,
std::hash<tvm::tir::VarNode const*>, std::equal_to<tvm::tir::VarNode const*>,
std::allocator<std::pair<tvm::tir::VarNode const* const, tvm::arith::IntSet> >
> const&, tvm::arith::Analyzer*)
arith::EstimateRegionUpperBound(region, var_dom, predicate, analyzer);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 1181, in
tvm::arith::EstimateRegionUpperBound(tvm::runtime::Array<tvm::Range, void>
const&, tvm::runtime::Map<tvm::tir::Var, tvm::Range, void, void> const&,
tvm::PrimExpr const&, tvm::arith::Analyzer*)
result.push_back(EvalSet(range, AsIntSet(var_dom)));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 1001, in
tvm::arith::EvalSet(tvm::Range, tvm::runtime::Map<tvm::tir::Var,
tvm::arith::IntSet, void, void> const&)
auto res = m.Eval(IntervalSet(r->min, ana.Simplify(sum)));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 379, in
tvm::arith::IntervalSetEvaluator::Eval(tvm::arith::IntervalSet)
IntervalSet min_set = this->Eval(val->min_value);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 373, in
tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
IntervalSet Eval(const PrimExpr& val) { return this->VisitExpr(val); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 428, in
tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::AddNode const*)
IntervalSet VisitExpr_(const AddNode* op) final { return
VisitBinaryExpr_<Add>(op); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 555, in
tvm::arith::IntervalSet
tvm::arith::IntervalSetEvaluator::VisitBinaryExpr_<tvm::tir::Add,
tvm::tir::AddNode>(tvm::tir::AddNode const*)
IntervalSet a = this->Eval(op->a);
^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 373, in
tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
IntervalSet Eval(const PrimExpr& val) { return this->VisitExpr(val); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 428, in
tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::AddNode const*)
IntervalSet VisitExpr_(const AddNode* op) final { return
VisitBinaryExpr_<Add>(op); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 555, in
tvm::arith::IntervalSet
tvm::arith::IntervalSetEvaluator::VisitBinaryExpr_<tvm::tir::Add,
tvm::tir::AddNode>(tvm::tir::AddNode const*)
IntervalSet a = this->Eval(op->a);
^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 373, in
tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
IntervalSet Eval(const PrimExpr& val) { return this->VisitExpr(val); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 428, in
tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::AddNode const*)
IntervalSet VisitExpr_(const AddNode* op) final { return
VisitBinaryExpr_<Add>(op); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 555, in
tvm::arith::IntervalSet
tvm::arith::IntervalSetEvaluator::VisitBinaryExpr_<tvm::tir::Add,
tvm::tir::AddNode>(tvm::tir::AddNode const*)
IntervalSet a = this->Eval(op->a);
^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 373, in
tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
IntervalSet Eval(const PrimExpr& val) { return this->VisitExpr(val); }
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 493, in
tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::BroadcastNodeconst*)
ICHECK(eval_vec_);
^^^^^
tvm.error.InternalError: Traceback (most recent call last):
30: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule,
tvm::transform::PassContext const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/transform.cc:121
29: operator()
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:755
28: tvm::tir::CompactBufferAllocation(tvm::tir::PrimFunc, bool)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:745
27: tvm::tir::BufferAccessRegionCollector::Collect(tvm::tir::PrimFunc
const&, bool)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:114
26:
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::AllocateNode const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:314
25:
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::AllocateNode const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:314
24:
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::AllocateNode const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:314
23: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:119
22: 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
21: operator()
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:119
20: tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::ForNode
const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:159
19: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:48
18:
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::LetStmtNode const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:170
17:
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::LetStmtNode const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:170
16:
tvm::tir::BufferAccessRegionCollector::VisitStmt_(tvm::tir::BufferStoreNode
const*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:134
15:
tvm::tir::BufferAccessRegionCollector::VisitBufferAccess(tvm::tir::BufferRegion
const&)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:377
14: tvm::tir::NDIntSetEval(tvm::runtime::Array<tvm::Range, void>,
tvm::PrimExpr, std::unordered_map<tvm::tir::VarNode const*, tvm::arith::IntSet,
std::hash<tvm::tir::VarNode const*>, std::equal_to<tvm::tir::VarNode const*>,
std::allocator<std::pair<tvm::tir::VarNode const* const, tvm::arith::IntSet> >
> const&, tvm::arith::Analyzer*)
at
/data/qshenaf/envs/tvm/src/tir/transforms/compact_buffer_region.cc:54
13: tvm::arith::EstimateRegionUpperBound(tvm::runtime::Array<tvm::Range,
void> const&, tvm::runtime::Map<tvm::tir::Var, tvm::Range, void, void> const&,
tvm::PrimExpr const&, tvm::arith::Analyzer*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:1181
12: tvm::arith::EvalSet(tvm::Range, tvm::runtime::Map<tvm::tir::Var,
tvm::arith::IntSet, void, void> const&)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:1001
11: tvm::arith::IntervalSetEvaluator::Eval(tvm::arith::IntervalSet)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:379
10: tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:373
9: tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::AddNode const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:428
8: tvm::arith::IntervalSet
tvm::arith::IntervalSetEvaluator::VisitBinaryExpr_<tvm::tir::Add,
tvm::tir::AddNode>(tvm::tir::AddNode const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:555
7: tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:373
6: tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::AddNode const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:428
5: tvm::arith::IntervalSet
tvm::arith::IntervalSetEvaluator::VisitBinaryExpr_<tvm::tir::Add,
tvm::tir::AddNode>(tvm::tir::AddNode const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:555
4: tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:373
3: tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::AddNode const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:428
2: tvm::arith::IntervalSet
tvm::arith::IntervalSetEvaluator::VisitBinaryExpr_<tvm::tir::Add,
tvm::tir::AddNode>(tvm::tir::AddNode const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:555
1: tvm::arith::IntervalSetEvaluator::Eval(tvm::PrimExpr const&)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:373
0: tvm::arith::IntervalSetEvaluator::VisitExpr_(tvm::tir::BroadcastNode
const*)
at /data/qshenaf/envs/tvm/src/arith/int_set.cc:493
File "/data/qshenaf/envs/tvm/src/arith/int_set.cc", line 493
InternalError: Check failed: (eval_vec_) is false:
```
### Environment
tvm-0.21.dev0
### Steps to reproduce
```
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(mod=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.ConvertBlocksToOpaque(),
tir.transform.FlattenBuffer(),tir.transform.NarrowDataType(32),tir.transform.LoopPartition(),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') # crash here
```
### Triage
* needs-triage
* 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]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
