Hi All, I’m writing an OP with TVMScript for a special backend, which has special data alignment requirements like this:
import tvm
from tvm.script import parser as I, tir as T
@T.prim_func
def test():
T.func_attr({"global_symbol": "main", "tir.noalias": True})
A = T.alloc_buffer((2, 8, 16), "float32", strides=(200, 16, 1), scope="local")
for ax0, ax1, ax2 in T.grid(2, 8, 16):
with T.block(""):
T.reads()
T.writes(A[ax0, ax1, ax2])
T.block_attr({"buffer_dim_align": [[0, 0, 200, 0]]})
A[ax0, ax1, ax2] = 4
mod = tvm.IRModule.from_expr(test)
res = tvm.tir.transform.LowerOpaqueBlock()(mod)
but an error occured while doing LowerOpaqueBlock:
Traceback (most recent call last):
File "test.py", line 42, in <module>
res = tvm.tir.transform.LowerOpaqueBlock()(mod)
File "/tvm/python/tvm/ir/transform.py", line 160, in __call__
return _ffi_transform_api.RunPass(self, mod)
File "/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line 237, in __call__
raise get_last_ffi_error()
tvm._ffi.base.TVMError: Traceback (most recent call last):
31: TVMFuncCall
at /tvm/src/runtime/c_runtime_api.cc:477
30: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const
at /tvm/include/tvm/runtime/packed_func.h:1217
29: Call
at /tvm/include/tvm/runtime/packed_func.h:1213
28: operator()
at /tvm/include/tvm/runtime/packed_func.h:1734
27: unpack_call<tvm::IRModule, 2, tvm::transform::<lambda(tvm::transform::Pass, tvm::IRModule)> >
at /tvm/include/tvm/runtime/packed_func.h:1674
26: run<>
at /tvm/include/tvm/runtime/packed_func.h:1634
25: run<tvm::runtime::TVMMovableArgValueWithContext_>
at /tvm/include/tvm/runtime/packed_func.h:1634
24: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_>
at /tvm/include/tvm/runtime/packed_func.h:1649
23: operator()
at /tvm/src/ir/transform.cc:499
22: tvm::transform::Pass::operator()(tvm::IRModule) const
at /tvm/src/ir/transform.cc:258
21: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
at /tvm/src/ir/transform.cc:274
20: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
at /tvm/src/tir/ir/transform.cc:101
19: tvm::runtime::TypedPackedFunc<tvm::tir::PrimFunc (tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)>::operator()(tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext) const
at /tvm/include/tvm/runtime/packed_func.h:1753
18: tvm::tir::PrimFunc tvm::runtime::detail::typed_packed_call_dispatcher<tvm::tir::PrimFunc>::run<tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext>(tvm::runtime::PackedFunc const&, tvm::tir::PrimFunc&&, tvm::IRModule&&, tvm::transform::PassContext&&)
at /tvm/include/tvm/runtime/packed_func.h:1697
17: tvm::runtime::TVMRetValue tvm::runtime::PackedFunc::operator()<tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext>(tvm::tir::PrimFunc&&, tvm::IRModule&&, tvm::transform::PassContext&&) const
at /tvm/include/tvm/runtime/packed_func.h:1621
16: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const
at /tvm/include/tvm/runtime/packed_func.h:1217
15: Call
at /tvm/include/tvm/runtime/packed_func.h:1213
14: operator()
at /tvm/include/tvm/runtime/packed_func.h:1747
13: unpack_call<tvm::tir::PrimFunc, 3, tvm::tir::transform::LowerOpaqueBlock()::<lambda(tvm::tir::PrimFunc, tvm::IRModule, tvm::transform::PassContext)> >
at /tvm/include/tvm/runtime/packed_func.h:1674
12: run<>
at /tvm/include/tvm/runtime/packed_func.h:1634
11: run<tvm::runtime::TVMMovableArgValueWithContext_>
at /tvm/include/tvm/runtime/packed_func.h:1634
10: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_>
at /tvm/include/tvm/runtime/packed_func.h:1634
9: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_>
at /tvm/include/tvm/runtime/packed_func.h:1649
8: operator()
at /tvm/src/tir/transforms/lower_opaque_block.cc:206
7: tvm::tir::LowerOpaqueBlock(tvm::tir::PrimFunc)
at /tvm/src/tir/transforms/lower_opaque_block.cc:195
6: tvm::tir::StmtMutator::operator()(tvm::tir::Stmt)
at /tvm/include/tvm/tir/stmt_functor.h:189
5: tvm::tir::StmtMutator::VisitStmt(tvm::tir::Stmt const&)
at /tvm/include/tvm/tir/stmt_functor.h:239
4: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::VisitStmt(tvm::tir::Stmt const&)
at /tvm/include/tvm/tir/stmt_functor.h:82
3: tvm::NodeFunctor<tvm::tir::Stmt (tvm::runtime::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)>::operator()(tvm::runtime::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
at /tvm/include/tvm/node/functor.h:97
2: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::runtime::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#18}::_FUN(tvm::runtime::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)
at /tvm/include/tvm/tir/stmt_functor.h:128
1: tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>::InitVTable()::{lambda(tvm::runtime::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*)#18}::operator()(tvm::runtime::ObjectRef const&, tvm::tir::StmtFunctor<tvm::tir::Stmt (tvm::tir::Stmt const&)>*) const
at /tvm/include/tvm/tir/stmt_functor.h:128
0: tvm::tir::OpaqueBlockLower::VisitStmt_(tvm::tir::BlockRealizeNode const*)
at /tvm/src/tir/transforms/lower_opaque_block.cc:56
File "/tvm/src/tir/transforms/lower_opaque_block.cc", line 56
TVMError:
---------------------------------------------------------------
An error occurred during the execution of TVM.
For more information, please see: https://tvm.apache.org/docs/errors.html
---------------------------------------------------------------
Check failed: (is_zero(floormod(buffer->strides[i - 1], buffer->strides[i]))) is false:
I noticed that this error is caused by 200%16 != 0, but I don’t know why? or how should I write an OP with special data alignment requirements?