Asuka0630 opened a new issue, #18472:
URL: https://github.com/apache/tvm/issues/18472
### Description
Using the meta schedule forces the use of the MMA version of
MultiLevelTilingTensorCore, which cannot generate valid valid code.
### Expected behavior
use of PTX MMA instructions
### Actual behavior
A fatal error occurred
### Environment
Python 3.12.7
TVM 0.23.0dev
### Steps to reproduce
``` python
import numpy as np
from typing import Tuple
import tvm
from tvm import te
from tvm import meta_schedule as ms
from tvm.te import create_prim_func
from tvm.meta_schedule.runner.config import EvaluatorConfig
from tvm.meta_schedule.space_generator import PostOrderApply
from tvm.meta_schedule.schedule_rule import ScheduleRule
import tvm.tir.tensor_intrin # pylint: disable=unused-import
M, N, K = 1024, 1024, 1024
def matmul(
m: int, n: int, k: int, in_dtype: str = "float16", out_dtype: str =
"float16"
) -> Tuple[te.Tensor, te.Tensor, te.Tensor]:
a = te.placeholder((m, k), name="A", dtype=in_dtype)
b = te.placeholder((k, n), name="B", dtype=in_dtype)
k = te.reduce_axis((0, k), name="k")
c = te.compute(
(m, n),
lambda i, j: te.sum(
a[i, k].astype(out_dtype) * b[k, j].astype(out_dtype), axis=[k]
),
name="C",
)
return (a, b, c)
if __name__ == "__main__":
np.random.seed(0)
target = tvm.target.Target.from_device(tvm.cuda(0))
sch_rules = ScheduleRule.create("cuda-tensorcore")
sch_rules = list(sch_rules)
sch_rules.pop(1)
space_gen = PostOrderApply(
sch_rules=sch_rules,
)
db = ms.tune_tir(
mod=create_prim_func(matmul(M, N, K)),
target=target,
max_trials_global=320,
num_trials_per_iter=16,
work_dir="./mma_gemm_test_fp16",
runner=ms.runner.LocalRunner(
evaluator_config=EvaluatorConfig(
number=3, min_repeat_ms=100, enable_cpu_cache_flush=False
),
),
cost_model=ms.cost_model.XGBModel(
extractor=ms.feature_extractor.PerStoreFeature(),
adaptive_training=False,
),
strategy=ms.search_strategy.EvolutionarySearch(),
special_space={"main": space_gen},
)
```
### Error Log
```
Traceback (most recent call last):
File "/home/tvm/python/tvm/exec/popen_worker.py", line 87, in main
result = fn(*args, **kwargs)
^^^^^^^^^^^^^^^^^^^
File "/home/tvm/python/tvm/meta_schedule/builder/local_builder.py", line
165, in <lambda>
lambda x: _worker_func(*x),
^^^^^^^^^^^^^^^^
File "/home/tvm/python/tvm/meta_schedule/builder/local_builder.py", line
231, in _worker_func
rt_mod: Module = f_build(mod, target, _deserialize_params(params))
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "python/tvm_ffi/cython/function.pxi", line 758, in
core.Function.__call__
File "python/tvm_ffi/cython/function.pxi", line 914, in
core.tvm_ffi_callback
File "/home/tvm/python/tvm/meta_schedule/builder/local_builder.py", line
261, in default_build
return tvm_build(mod, target=target)
File "/home/tvm/python/tvm/driver/build_module.py", line 59, in build
return tvm.tir.build(mod, target, pipeline)
File "/home/tvm/python/tvm/tir/build.py", line 226, in build
mod = pipeline(mod)
File "/home/tvm/python/tvm/ir/transform.py", line 167, in __call__
return _ffi_transform_api.RunPass(self, mod)
File "python/tvm_ffi/cython/function.pxi", line 758, in
core.Function.__call__
File "/home/tvm/src/ir/transform.cc", line 550, in
unpack_call<tvm::IRModule, 0, 1,
tvm::transform::__TVMFFIStaticInitFunc4()::<lambda(tvm::transform::Pass,
tvm::ffi::RValueRef<tvm::IRModule>)> >
[](Pass pass, ffi::RValueRef<IRModule> mod) { return
pass(*std::move(mod)); });
File "/home/tvm/src/ir/transform.cc", line 296, in
tvm::transform::Pass::operator()(tvm::IRModule) const
return this->operator()(std::move(mod), PassContext::Current());
File "/home/tvm/src/ir/transform.cc", line 312, in
tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext
const&) const
ret = node->operator()(std::move(mod), pass_ctx);
File "/home/tvm/src/ir/transform.cc", line 420, in
tvm::transform::ModulePassNode::operator()(tvm::IRModule,
tvm::transform::PassContext const&) const
mod = pass_func(std::move(mod), pass_ctx);
File "/home/tvm/src/ir/transform.cc", line 545, in operator()
return pass_func(ffi::RValueRef<IRModule>(std::move(mod)), ctx);
File "python/tvm_ffi/cython/function.pxi", line 914, in
core.tvm_ffi_callback
File "/home/tvm/python/tvm/tir/pipeline.py", line 123, in _pipeline
mod = tvm.ir.transform.Sequential(passes)(mod)
File "/home/tvm/python/tvm/ir/transform.py", line 167, in __call__
return _ffi_transform_api.RunPass(self, mod)
File "python/tvm_ffi/cython/function.pxi", line 758, in
core.Function.__call__
File "/home/tvm/src/ir/transform.cc", line 550, in
unpack_call<tvm::IRModule, 0, 1,
tvm::transform::__TVMFFIStaticInitFunc4()::<lambda(tvm::transform::Pass,
tvm::ffi::RValueRef<tvm::IRModule>)> >
[](Pass pass, ffi::RValueRef<IRModule> mod) { return
pass(*std::move(mod)); });
File "/home/tvm/src/ir/transform.cc", line 296, in
tvm::transform::Pass::operator()(tvm::IRModule) const
return this->operator()(std::move(mod), PassContext::Current());
File "/home/tvm/src/ir/transform.cc", line 312, in
tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext
const&) const
ret = node->operator()(std::move(mod), pass_ctx);
File "/home/tvm/src/ir/transform.cc", line 490, in
tvm::transform::SequentialNode::operator()(tvm::IRModule,
tvm::transform::PassContext const&) const
mod = pass(std::move(mod), pass_ctx);
File "/home/tvm/src/ir/transform.cc", line 312, in
tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext
const&) const
ret = node->operator()(std::move(mod), pass_ctx);
File "/home/tvm/src/tir/ir/transform.cc", line 124, in
tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule,
tvm::transform::PassContext const&) const
func = pass_func(std::move(func), mod, pass_ctx);
File "/home/tvm/src/tir/transforms/transform_mma_buffer_layout.cc", line
184, in operator()
n->body = MmaBufferLayoutTransformer()(std::move(n->body));
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 182, in
tvm::tir::StmtMutator::operator()(tvm::tir::Stmt)
return VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 123, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#15}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BlockRealizeNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 123, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#15}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BlockRealizeNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 489, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::BlockRealizeNode const*)
Stmt block = this->VisitStmt(op->block);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 122, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#14}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BlockNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 122, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#14}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BlockNode);
File "/home/tvm/src/tir/transforms/transform_mma_buffer_layout.cc", line
124, in tvm::tir::MmaBufferLayoutTransformer::VisitStmt_(tvm::tir::BlockNode
const*)
n->body = VisitStmt(n->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 232, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
Stmt ret = StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 123, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#15}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BlockRealizeNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 123, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#15}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BlockRealizeNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 489, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::BlockRealizeNode const*)
Stmt block = this->VisitStmt(op->block);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 122, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#14}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BlockNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 122, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#14}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BlockNode);
File "/home/tvm/src/tir/transforms/transform_mma_buffer_layout.cc", line
124, in tvm::tir::MmaBufferLayoutTransformer::VisitStmt_(tvm::tir::BlockNode
const*)
n->body = VisitStmt(n->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 118, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#10}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(SeqStmtNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 118, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#10}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(SeqStmtNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 380, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::SeqStmtNode const*)
ffi::Array<Stmt> seq = Internal::Mutate(this, op->seq);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 191, in
tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)
return MutateArray(self, arr, fmutate);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 164, in
tvm::ffi::Array<tvm::tir::Stmt,
std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type>
tvm::tir::StmtMutator::Internal::MutateArray<tvm::tir::Stmt,
tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt
const&)#1}>(tvm::tir::StmtMutator*, tvm::ffi::Array<tvm::tir::Stmt,
std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type> const&,
tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt
const&)#1})
ffi::Array<T> copy = arr.Map(fmutate);
File "/home/tvm/3rdparty/tvm-ffi/include/tvm/ffi/container/array.h", line
799, in tvm::ffi::Array<tvm::tir::Stmt,
std::enable_if<storage_enabled_v<tvm::tir::Stmt>, void>::type>
tvm::ffi::Array<tvm::tir::Stmt,
void>::Map<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt
const&)#1},
tvm::tir::Stmt>(tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt
const&)#1}) const
return Array<U>(MapHelper(data_, fmap));
File "/home/tvm/3rdparty/tvm-ffi/include/tvm/ffi/container/array.h", line
975, in tvm::ffi::ObjectPtr<tvm::ffi::Object> tvm::ffi::Array<tvm::tir::Stmt,
void>::MapHelper<tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt
const&)#1}, tvm::tir::Stmt>(tvm::ffi::ObjectPtr<tvm::ffi::Object>,
tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::Array<tvm::tir::Stmt, void> const&)::{lambda(tvm::tir::Stmt
const&)#1})
U mapped = fmap(details::AnyUnsafe::CopyFromAnyViewAfterCheck<T>(*it));
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 190, in
tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*,
tvm::ffi::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 "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 123, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#15}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BlockRealizeNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 123, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#15}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BlockRealizeNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 489, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::BlockRealizeNode const*)
Stmt block = this->VisitStmt(op->block);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 122, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#14}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BlockNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 122, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#14}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BlockNode);
File "/home/tvm/src/tir/transforms/transform_mma_buffer_layout.cc", line
124, in tvm::tir::MmaBufferLayoutTransformer::VisitStmt_(tvm::tir::BlockNode
const*)
n->body = VisitStmt(n->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 112, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#4}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(ForNode);
File "/home/tvm/src/tir/ir/stmt_functor.cc", line 263, in
tvm::tir::StmtMutator::VisitStmt_(tvm::tir::ForNode const*)
Stmt body = this->VisitStmt(op->body);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 236, in
tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
return StmtFunctor::VisitStmt(stmt);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 82, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::VisitStmt(tvm::tir::Stmt const&)
return vtable(n, this, std::forward<Args>(args)...);
File "/home/tvm/include/tvm/node/functor.h", line 102, in
tvm::NodeFunctor<tvm::tir::Stmt (tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)>::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
return (*func_[n->type_index() - begin_type_index_])(n,
std::forward<Args>(args)...);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 120, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#12}::_FUN(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
IR_STMT_FUNCTOR_DISPATCH(BufferStoreNode);
File "/home/tvm/include/tvm/tir/stmt_functor.h", line 120, in
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>::InitVTable()::{lambda(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt
const&)>*)#12}::operator()(tvm::ffi::ObjectRef const&,
tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
IR_STMT_FUNCTOR_DISPATCH(BufferStoreNode);
File "/home/tvm/src/tir/transforms/transform_mma_buffer_layout.cc", line
134, in
tvm::tir::MmaBufferLayoutTransformer::VisitStmt_(tvm::tir::BufferStoreNode
const*)
ICHECK(index_map_func.has_value());
File "/home/tvm/include/tvm/runtime/logging.h", line 321, in
tvm::runtime::detail::LogFatal::~LogFatal()
GetEntry().Finalize();
File "/home/tvm/include/tvm/runtime/logging.h", line 337, in
tvm::runtime::detail::LogFatal::Entry::Finalize()
InternalError error(file_, lineno_, stream_.str());
tvm.error.InternalError: Check failed: (index_map_func.has_value()) is
false:
```
### Triage
* needs-triage
* bug
--
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]
