Skip to content

[Bug] InternalError: Check failed: (eval_vec_) is false #17936

Open
Open
[Bug] InternalError: Check failed: (eval_vec_) is false#17936
Labels
needs-triagePRs or issues that need to be investigated by maintainers to find the right assignees to address ittype: bug
@Cookiee235

Description

@Cookiee235
Cookiee235
opened on May 9, 2025

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

Metadata

Metadata

Assignees

No one assigned

    Labels

    needs-triagePRs or issues that need to be investigated by maintainers to find the right assignees to address ittype: bug

    Type

    No type

    Fields

    No fields configured for issues without a type.

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions