Apache TVM 是一個端到端的深度學習編譯框架，適用於 CPU、GPU 和各種機器學習加速芯片。更多 TVM 中文文檔可訪問 →https://tvm.hyper.ai/
作者：Lianmin Zheng
本教程介紹如何為 NVIDIA GPU 編寫高性能可調模板。通過在此模板上運行自動調優器，可在許多情況下勝過供應商提供的 cuDNN 庫。
注意，本教程不會在 Windows 或最新版本的 macOS 上運行。如需運行，請將本教程的主體放在 if name == "__main__": 代碼塊中。
安裝依賴
要在 TVM 中使用 autotvm 包，需要安裝額外的依賴（如果用的是 Python2，請將「3」更改為「2」）：
pip3 install --user psutil xgboost tornado cloudpickle為了讓 TVM 在調優中運行更快，推薦使用 Cython 作為 TVM 的 FFI。在 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。
但是，它們的實現是針對一些特殊的輸入 shape 手動調整的。本節將構建足夠大的空間，涵蓋這些教程中使用的技術，然後依靠高效的自動調優器，對空間進行搜索並選擇合適的配置。
熟悉 CUDA schedule 的開發者，對以下通用模板並不陌生。可以修改此模板以調整其他算子，例如深度卷積和 GEMM。要完全理解這個模板，需要熟悉調度原語和自動調優 API，可以參考上面的教程和 AutoTVM 教程。
需要注意 conv2d 算子的搜索空間可能非常大（某些輸入 shape 為 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
