Apache TVMは、CPU、GPU、そして様々な機械学習アクセラレータチップに適した、エンドツーエンドのディープラーニング構築フレームワークです。中国語版のTVMドキュメントは、→ https://tvm.hyper.ai/ をご覧ください。

著者: 鄭連民

このチュートリアルでは、NVIDIA GPU向けに高性能かつチューニング可能なテンプレートを作成する方法について説明します。このテンプレートに自動チューナーを実行することで、多くの場合、ベンダー提供のcuDNNライブラリよりも優れたパフォーマンスを実現できます。

このチュートリアルはWindowsまたは最新バージョンのmacOSでは実行できませんのでご注意ください。実行するには、チュートリアルの本文を `if name == "__main__":` コードブロック内に配置してください。

依存関係をインストールする

TVM で autotvm パッケージを使用するには、追加の依存関係をインストールする必要があります (Python 2 を使用している場合は、「3」を「2」に変更してください)。

 pip3 install --user psutil xgboost tornado cloudpickle

チューニング中のTVMパフォーマンスを向上させるには、TVMのFFIとしてCythonを使用することをお勧めします。TVMのルートディレクトリで、以下のコマンドを実行してください。

 pip3 install --user cython sudo make cython3

Python コードでパッケージをインポートします。

 import logging import sys import numpy as np import tvm from tvm import te, topi, testing from tvm.topi.testing import conv2d_nchw_python import tvm.testing from tvm import autotvm

ステップ1: 検索空間を定義する

TVMには、(1)GPUで畳み込みを最適化する方法や、(2)NVIDIA GPUでDepthwiseConvを最適化する方法など、多くの便利なスケジューリングプリミティブと詳細なチュートリアルが含まれています。

しかし、これらの実装は特定の入力形状に合わせて手動で調整されています。このセクションでは、これらのチュートリアルで使用される手法をカバーするのに十分な大きさの空間を構築し、効率的な自動チューナーを使用して空間を探索し、適切な構成を選択します。

CUDAスケジュールに精通している開発者であれば、以下の一般的なテンプレートはお分かりいただけるでしょう。このテンプレートは、深度方向畳み込みやGEMMなどの他の演算子に合わせて変更できます。このテンプレートを完全に理解するには、スケジューリングプリミティブとAutoTVM APIに関する知識が必要です。上記のチュートリアルとAutoTVMチュートリアルを参照してください。

conv2d 演算子の検索空間は非常に大きくなる可能性があることに注意してください (一部の入力形状は 10^9 のオーダーになります)。

 @autotvm.template("tutorial/conv2d_no_batching") def conv2d_no_batching(N, H, W, CO, CI, KH, KW, stride, padding): assert N == 1, "Only consider batch_size = 1 in this template" data = te.placeholder((N, CI, H, W), name="data") kernel = te.placeholder((CO, CI, KH, KW), name="kernel") conv = topi.nn.conv2d_nchw(data, kernel, stride, padding, dilation=1, out_dtype="float32") s = te.create_schedule([conv.op]) #### 空间定义开始#### n, f, y, x = s[conv].op.axis rc, ry, rx = s[conv].op.reduce_axis cfg = autotvm.get_config() cfg.define_split("tile_f", f, num_outputs=4) cfg.define_split("tile_y", y, num_outputs=4) cfg.define_split("tile_x", x, num_outputs=4) cfg.define_split("tile_rc", rc, num_outputs=3) cfg.define_split("tile_ry", ry, num_outputs=3) cfg.define_split("tile_rx", rx, num_outputs=3) cfg.define_knob("auto_unroll_max_step", [0, 512, 1500]) cfg.define_knob("unroll_explicit", [0, 1]) #### 空间定义结束#### # 内联填充pad_data = s[conv].op.input_tensors[0] s[pad_data].compute_inline() data, raw_data = pad_data, data output = conv OL = s.cache_write(conv, "local") # 创建cache 阶段AA = s.cache_read(data, "shared", [OL]) WW = s.cache_read(kernel, "shared", [OL]) AL = s.cache_read(AA, "local", [OL]) WL = s.cache_read(WW, "local", [OL]) # 平铺和绑定空间轴n, f, y, x = s[output].op.axis bf, vf, tf, fi = cfg["tile_f"].apply(s, output, f) by, vy, ty, yi = cfg["tile_y"].apply(s, output, y) bx, vx, tx, xi = cfg["tile_x"].apply(s, output, x) kernel_scope = n # 这是在此内核中附加全局配置的范围s[output].bind(bf, te.thread_axis("blockIdx.z")) s[output].bind(by, te.thread_axis("blockIdx.y")) s[output].bind(bx, te.thread_axis("blockIdx.x")) s[output].bind(vf, te.thread_axis("vthread")) s[output].bind(vy, te.thread_axis("vthread")) s[output].bind(vx, te.thread_axis("vthread")) s[output].bind(tf, te.thread_axis("threadIdx.z")) s[output].bind(ty, te.thread_axis("threadIdx.y")) s[output].bind(tx, te.thread_axis("threadIdx.x")) s[output].reorder(n, bf, by, bx, vf, vy, vx, tf, ty, tx, fi, yi, xi) s[OL].compute_at(s[output], tx) # tile reduction 轴n, f, y, x = s[OL].op.axis rc, ry, rx = s[OL].op.reduce_axis rco, rcm, rci = cfg["tile_rc"].apply(s, OL, rc) ryo, rym, ryi = cfg["tile_rx"].apply(s, OL, ry) rxo, rxm, rxi = cfg["tile_ry"].apply(s, OL, rx) s[OL].reorder(rco, ryo, rxo, rcm, rym, rxm, rci, ryi, rxi, n, f, y, x) s[AA].compute_at(s[OL], rxo) s[WW].compute_at(s[OL], rxo) s[AL].compute_at(s[OL], rxm) s[WL].compute_at(s[OL], rxm) # 协作获取for load in [AA, WW]: n, f, y, x = s[load].op.axis fused = s[load].fuse(n, f, y, x) tz, fused = s[load].split(fused, nparts=cfg["tile_f"].size[2]) ty, fused = s[load].split(fused, nparts=cfg["tile_y"].size[2]) tx, fused = s[load].split(fused, nparts=cfg["tile_x"].size[2]) s[load].bind(tz, te.thread_axis("threadIdx.z")) s[load].bind(ty, te.thread_axis("threadIdx.y")) s[load].bind(tx, te.thread_axis("threadIdx.x")) # 调优unroll s[output].pragma(kernel_scope, "auto_unroll_max_step", cfg["auto_unroll_max_step"].val) s[output].pragma(kernel_scope, "unroll_explicit", cfg["unroll_explicit"].val) return s, [raw_data, kernel, conv]

ステップ2: 宇宙での検索

ResNetの最終層をテストケースとして選択しました。十分なメモリ容量を持つXGBoostTunerが最適です。ここではデモ用に20回の試行のみを実行しています。実際には1000回の試行で、このテンプレートに適したカーネルを見つけることができるでしょう。

 # logging 配置（用于将调优日志打印到屏幕） logging.getLogger("autotvm").setLevel(logging.DEBUG) logging.getLogger("autotvm").addHandler(logging.StreamHandler(sys.stdout)) # resnet 中的最后一层N, H, W, CO, CI, KH, KW, strides, padding = 1, 7, 7, 512, 512, 3, 3, (1, 1), (1, 1) task = autotvm.task.create( "tutorial/conv2d_no_batching", args=(N, H, W, CO, CI, KH, KW, strides, padding), target="cuda" ) print(task.config_space) # 使用本地gpu，为每个配置测量10 次以减少方差# 编译程序超时10 秒，运行超时4 秒measure_option = autotvm.measure_option( builder=autotvm.LocalBuilder(), runner=autotvm.LocalRunner(repeat=3, min_repeat_ms=100, timeout=4), ) record_file = None # 开始调优，将log 记录到`conv2d.log` # 在调优过程中，会尝试很多无效的配置，所以你应该# 查看许多错误报告。只要能看到非零的GFLOPS 就可以。 # 我们不再在服务器上运行调优，因为太耗时间了# 去掉下面的注释，自己运行# tuner = autotvm.tuner.XGBTuner(task) # record_file = "conv2d.log" # tuner.tune( # n_trial=5, # measure_option=measure_option, # callbacks=[autotvm.callback.log_to_file(record_file)], # )

出力結果:

 ConfigSpace (len=10454400, space_map= 0 tile_f: Split(policy=factors, product=512, num_outputs=4) len=220 1 tile_y: Split(policy=factors, product=7, num_outputs=4) len=4 2 tile_x: Split(policy=factors, product=7, num_outputs=4) len=4 3 tile_rc: Split(policy=factors, product=512, num_outputs=3) len=55 4 tile_ry: Split(policy=factors, product=3, num_outputs=3) len=3 5 tile_rx: Split(policy=factors, product=3, num_outputs=3) len=3 6 auto_unroll_max_step: OtherOption([0, 512, 1500]) len=3 7 unroll_explicit: OtherOption([0, 1]) len=2 ) waiting for device... device available Get devices for measurement successfully! No: 1 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 32, 8, 2]), ('tile_y', [-1, 1, 1, 1]), ('tile_x', [-1, 7, 1, 1]), ('tile_rc', [-1, 8, 64]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 3, 1]), ('auto_unroll_max_step', 0), ('unroll_explicit', 1)],None,6171524 No: 2 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 128, 1, 4]), ('tile_y', [-1, 1, 1, 1]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 4, 128]), ('tile_ry', [-1, 1, 1]), ('tile_rx', [-1, 3, 1]), ('auto_unroll_max_step', 512), ('unroll_explicit', 0)],None,2502827 No: 3 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 4, 1, 32]), ('tile_y', [-1, 1, 1, 1]), ('tile_x', [-1, 1, 1, 7]), ('tile_rc', [-1, 512, 1]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 512), ('unroll_explicit', 0)],None,3325707 No: 4 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 8, 4, 2]), ('tile_y', [-1, 1, 1, 7]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 4, 8]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,4942815 No: 5 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 4, 1, 128]), ('tile_y', [-1, 1, 1, 7]), ('tile_x', [-1, 7, 1, 1]), ('tile_rc', [-1, 2, 64]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,5197272 No: 6 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 32, 2, 4]), ('tile_y', [-1, 1, 1, 1]), ('tile_x', [-1, 1, 7, 1]), ('tile_rc', [-1, 8, 8]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 1]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,3979473 No: 7 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 8, 4, 8]), ('tile_y', [-1, 1, 7, 1]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 4, 32]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 512), ('unroll_explicit', 0)],None,3439632 No: 8 GFLOPS: 0.00/0.00 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 142, in build res = future.result() File "/usr/lib/python3.7/concurrent/futures/_base.py", line 435, in result return self.__get_result() File "/usr/lib/python3.7/concurrent/futures/_base.py", line 384, in __get_result raise self._exception File "/usr/lib/python3.7/concurrent/futures/thread.py", line 57, in run result = self.fn(*self.args, **self.kwargs) File "/workspace/python/tvm/contrib/popen_pool.py", line 404, in <lambda> worker = lambda *args: self._worker_run(*args) File "/workspace/python/tvm/contrib/popen_pool.py", line 373, in _worker_run return proc.recv() File "/workspace/python/tvm/contrib/popen_pool.py", line 297, in recv raise TimeoutError() TimeoutError [('tile_f', [-1, 2, 1, 64]), ('tile_y', [-1, 1, 1, 7]), ('tile_x', [-1, 1, 7, 1]), ('tile_rc', [-1, 1, 4]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,4909501 No: 9 GFLOPS: 174.65/174.65 result: MeasureResult(costs=(0.0013254985555555556,), error_no=MeasureErrorNo.NO_ERROR, all_cost=2.0453152656555176, timestamp=1658801307.741073) [('tile_f', [-1, 1, 4, 8]), ('tile_y', [-1, 7, 1, 1]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 2, 2]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,5072689 No: 10 GFLOPS: 0.00/174.65 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 4, 4, 8]), ('tile_y', [-1, 1, 1, 1]), ('tile_x', [-1, 1, 1, 7]), ('tile_rc', [-1, 64, 2]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,5092711 No: 11 GFLOPS: 260.18/260.18 result: MeasureResult(costs=(0.0008897650828729283,), error_no=MeasureErrorNo.NO_ERROR, all_cost=1.7439014911651611, timestamp=1658801308.6655772) [('tile_f', [-1, 8, 2, 1]), ('tile_y', [-1, 7, 1, 1]), ('tile_x', [-1, 1, 7, 1]), ('tile_rc', [-1, 2, 1]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 3, 1]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,4264713 No: 12 GFLOPS: 0.00/260.18 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 128, 1, 2]), ('tile_y', [-1, 1, 7, 1]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 1, 256]), ('tile_ry', [-1, 1, 1]), ('tile_rx', [-1, 1, 1]), ('auto_unroll_max_step', 0), ('unroll_explicit', 0)],None,183542 No: 13 GFLOPS: 0.00/260.18 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 4, 8, 8]), ('tile_y', [-1, 1, 7, 1]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 1, 64]), ('tile_ry', [-1, 1, 1]), ('tile_rx', [-1, 3, 1]), ('auto_unroll_max_step', 512), ('unroll_explicit', 0)],None,2482196 No: 14 GFLOPS: 0.00/260.18 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 64, 1, 4]), ('tile_y', [-1, 1, 7, 1]), ('tile_x', [-1, 1, 1, 7]), ('tile_rc', [-1, 4, 2]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 1)],None,10306226 No: 15 GFLOPS: 5.42/260.18 result: MeasureResult(costs=(0.042678113000000004,), error_no=MeasureErrorNo.NO_ERROR, all_cost=1.8313236236572266, timestamp=1658801313.2114427) [('tile_f', [-1, 2, 2, 8]), ('tile_y', [-1, 1, 1, 7]), ('tile_x', [-1, 7, 1, 1]), ('tile_rc', [-1, 4, 8]), ('tile_ry', [-1, 1, 1]), ('tile_rx', [-1, 1, 1]), ('auto_unroll_max_step', 0), ('unroll_explicit', 1)],None,5330964 No: 16 GFLOPS: 3.34/260.18 result: MeasureResult(costs=(0.0693738225,), error_no=MeasureErrorNo.NO_ERROR, all_cost=4.546748638153076, timestamp=1658801314.4557946) [('tile_f', [-1, 8, 4, 4]), ('tile_y', [-1, 1, 1, 7]), ('tile_x', [-1, 1, 1, 7]), ('tile_rc', [-1, 4, 1]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 1]), ('auto_unroll_max_step', 512), ('unroll_explicit', 0)],None,2140058 No: 17 GFLOPS: 0.00/260.18 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 142, in build res = future.result() File "/usr/lib/python3.7/concurrent/futures/_base.py", line 435, in result return self.__get_result() File "/usr/lib/python3.7/concurrent/futures/_base.py", line 384, in __get_result raise self._exception File "/usr/lib/python3.7/concurrent/futures/thread.py", line 57, in run result = self.fn(*self.args, **self.kwargs) File "/workspace/python/tvm/contrib/popen_pool.py", line 404, in <lambda> worker = lambda *args: self._worker_run(*args) File "/workspace/python/tvm/contrib/popen_pool.py", line 373, in _worker_run return proc.recv() File "/workspace/python/tvm/contrib/popen_pool.py", line 297, in recv raise TimeoutError() TimeoutError [('tile_f', [-1, 2, 2, 1]), ('tile_y', [-1, 1, 7, 1]), ('tile_x', [-1, 7, 1, 1]), ('tile_rc', [-1, 4, 16]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 1)],None,10195251 No: 18 GFLOPS: 27.80/260.18 result: MeasureResult(costs=(0.008326810928571429,), error_no=MeasureErrorNo.NO_ERROR, all_cost=1.2733991146087646, timestamp=1658801325.4546382) [('tile_f', [-1, 4, 8, 4]), ('tile_y', [-1, 1, 1, 1]), ('tile_x', [-1, 1, 1, 1]), ('tile_rc', [-1, 1, 4]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 3, 1]), ('auto_unroll_max_step', 0), ('unroll_explicit', 1)],None,6068603 No: 19 GFLOPS: 0.00/260.18 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 16, 4, 8]), ('tile_y', [-1, 1, 7, 1]), ('tile_x', [-1, 7, 1, 1]), ('tile_rc', [-1, 4, 128]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 0), ('unroll_explicit', 1)],None,6956993 No: 20 GFLOPS: 0.00/260.18 result: Traceback (most recent call last): File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 588, in __call__ func, arg_info = _build_func_common(measure_input, self.runtime, **kwargs) File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 540, in _build_func_common func = build(s, args, target_host=task.target_host, runtime=runtime) File "/workspace/python/tvm/driver/build_module.py", line 228, in build input_mod = lower(inputs, args, name=name, binds=binds) File "/workspace/python/tvm/driver/build_module.py", line 134, in lower return ffi.lower_schedule(inp, args, name, binds, simple_mode) File "tvm/_ffi/_cython/./packed_func.pxi", line 331, in tvm._ffi._cy3.core.PackedFuncBase.__call__ File "tvm/_ffi/_cython/./packed_func.pxi", line 276, in tvm._ffi._cy3.core.FuncCall File "tvm/_ffi/_cython/./base.pxi", line 181, in tvm._ffi._cy3.core.CHECK_CALL tvm._ffi.base.TVMError: Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel Traceback (most recent call last): 24: TVMFuncCall at ../src/runtime/c_runtime_api.cc:477 23: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 22: Call at ../include/tvm/runtime/packed_func.h:1213 21: operator() at ../include/tvm/runtime/packed_func.h:1731 20: unpack_call<tvm::IRModule, 5, tvm::<lambda(tvm::te::Schedule, const tvm::runtime::Array<tvm::runtime::ObjectRef>&, const tvm::runtime::String&, const tvm::runtime::Map<tvm::te::Tensor, tvm::tir::Buffer>&, bool)> > at ../include/tvm/runtime/packed_func.h:1671 19: run<> at ../include/tvm/runtime/packed_func.h:1631 18: run<tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 17: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 16: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 15: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1631 14: run<tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_, tvm::runtime::TVMMovableArgValueWithContext_> at ../include/tvm/runtime/packed_func.h:1646 13: operator() at ../src/driver/driver_api.cc:365 12: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array<tvm::runtime::ObjectRef, void> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::unordered_map<tvm::te::Tensor, tvm::tir::Buffer, std::hash<tvm::te::Tensor>, std::equal_to<tvm::te::Tensor>, std::allocator<std::pair<tvm::te::Tensor const, tvm::tir::Buffer> > > const&, bool) at ../src/driver/driver_api.cc:352 11: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array<tvm::transform::Pass, void>) at ../src/driver/driver_api.cc:252 10: tvm::transform::Pass::operator()(tvm::IRModule) const at ../src/ir/transform.cc:258 9: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 8: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:453 7: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/ir/transform.cc:274 6: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const at ../src/tir/ir/transform.cc:100 5: 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 ../include/tvm/runtime/packed_func.h:1750 4: 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 ../include/tvm/runtime/packed_func.h:1694 3: 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 ../include/tvm/runtime/packed_func.h:1618 2: tvm::runtime::PackedFuncObj::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const at ../include/tvm/runtime/packed_func.h:1217 1: Call at ../include/tvm/runtime/packed_func.h:1213 0: operator() at ../src/runtime/c_runtime_api.cc:534 File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback File "/workspace/python/tvm/autotvm/measure/measure_methods.py", line 871, in verify_pass raise InstantiationError("Skipped because of invalid gpu kernel") tvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel [('tile_f', [-1, 16, 1, 2]), ('tile_y', [-1, 7, 1, 1]), ('tile_x', [-1, 1, 7, 1]), ('tile_rc', [-1, 32, 4]), ('tile_ry', [-1, 1, 3]), ('tile_rx', [-1, 1, 3]), ('auto_unroll_max_step', 512), ('unroll_explicit', 0)],None,3377719

最后从日志文件中检查最佳配置，检查正确性并测试运行时间。

 # 检查最佳配置dispatch_context = autotvm.apply_history_best(record_file) best_config = dispatch_context.query(task.target, task.workload) print("\nBest config:") print(best_config) # 从日志文件中应用历史最好记录with autotvm.apply_history_best(record_file): with tvm.target.Target("cuda"): s, arg_bufs = conv2d_no_batching(N, H, W, CO, CI, KH, KW, strides, padding) func = tvm.build(s, arg_bufs) # 验证正确性a_np = np.random.uniform(size=(N, CI, H, W)).astype(np.float32) w_np = np.random.uniform(size=(CO, CI, KH, KW)).astype(np.float32) c_np = conv2d_nchw_python(a_np, w_np, strides, padding) dev = tvm.cuda() a_tvm = tvm.nd.array(a_np, device=dev) w_tvm = tvm.nd.array(w_np, device=dev) c_tvm = tvm.nd.empty(c_np.shape, device=dev) func(a_tvm, w_tvm, c_tvm) tvm.testing.assert_allclose(c_np, c_tvm.numpy(), rtol=1e-2) # 评估运行时间。这里选择一个较大的重复次数（400）来减少噪音以及内核启动的开销。还可用nvprof 来验证结果。 evaluator = func.time_evaluator(func.entry_name, dev, number=400) print("Time cost of this operator: %f" % evaluator(a_tvm, w_tvm, c_tvm).mean)

出力結果:

 Finish loading 20 records Best config: [('tile_f', [-1, 8, 2, 1]), ('tile_y', [-1, 7, 1, 1]), ('tile_x', [-1, 1, 7, 1]), ('tile_rc', [-1, 2, 1]), ('tile_ry', [-1, 3, 1]), ('tile_rx', [-1, 3, 1]), ('auto_unroll_max_step', 1500), ('unroll_explicit', 0)],None,4264713 Finish loading 20 records Time cost of this operator: 0.001299

下载Python 源代码：tune_conv2d_cuda.py

下载Jupyter notebook：tune_conv2d_cuda.ipynb