[AutoTVM]Error for Tuning SSD Model on RK3399 Mali GPU

Compile...
WARNING:autotvm:Cannot find config for target=opencl -device=mali, workload=('conv2d_transpose_nchw', (1, 256, 5, 5, 'float32'), (256, 256, 2, 2, 'float32'), (2, 2), (0, 0), 'float32'). A fallback configuration is used, which may bring great performance regression.
[20:31:44] tvm/src/arithmetic/int_set.cc:256: Return Everything in CombineInterval Mul
....
[20:31:44] tvm/src/arithmetic/int_set.cc:282: Return Everything in CombineInterval Div
[20:31:44] tvm/src/arithmetic/int_set.cc:299: Return Everything in CombineInterval Mod
Traceback (most recent call last):
  File "tune_nnvm_mobile_gpu.py", line 362, in <module>
    tune_and_evaluate(tuning_option)
  File "tune_nnvm_mobile_gpu.py", line 325, in tune_and_evaluate
    shape={'data': input_shape}, params=params, dtype=dtype)
  File "tvm/nnvm/python/nnvm/compiler/build_module.py", line 305, in build
    graph = graph.apply("GraphCompile")
  File "tvm/nnvm/python/nnvm/graph.py", line 234, in apply
    check_call(_LIB.NNGraphApplyPasses(self.handle, npass, cpass, ctypes.byref(ghandle)))
  File "tvm/nnvm/python/nnvm/_base.py", line 75, in check_call
    raise NNVMError(py_str(_LIB.NNGetLastError()))
nnvm._base.NNVMError: TVMCall CFunc Error:
Traceback (most recent call last):
  File "tvm/_ffi/_cython/./function.pxi", line 38, in tvm._ffi._cy2.core.tvm_callback
    rv = local_pyfunc(*pyargs)
  File "tvm/nnvm/python/nnvm/compiler/build_module.py", line 115, in _lower
    raise RuntimeError(msg)
RuntimeError: Traceback (most recent call last):
  File "tvm/nnvm/python/nnvm/compiler/build_module.py", line 107, in _lower
    f = tvm.lower(sch, inputs, name=func_name)
  File "tvm/python/tvm/build_module.py", line 380, in lower
    return ir_pass.MakeAPI(stmt, name, arg_list, 0, cfg.restricted_func)
  File "tvm/_ffi/_cython/./function.pxi", line 286, in tvm._ffi._cy2.core.FunctionBase.__call__
    FuncCall(self.chandle, args, &ret_val, &ret_tcode)
  File "tvm/_ffi/_cython/./function.pxi", line 231, in tvm._ffi._cy2.core.FuncCall
    CALL(TVMFuncCall(chandle, &values[0], &tcodes[0],
  File "tvm/_ffi/_cython/./base.pxi", line 143, in tvm._ffi._cy2.core.CALL
    raise TVMError(py_str(TVMGetLastError()))
TVMError: [20:31:44] tvm/src/pass/make_api.cc:169: Not all Vars are passed in api_args:  'blockIdx.y'  'blockIdx.x'  does not appeared in api_args

Stack trace returned 10 entries:
[bt] (0) tvm/build/libtvm.so(dmlc::StackTrace[abi:cxx11]()+0x5b) [0x7f7e648258ab]
[bt] (1) tvm/build/libtvm.so(dmlc::LogMessageFatal::~LogMessageFatal()+0x28) [0x7f7e64826158]
[bt] (2) tvm/build/libtvm.so(tvm::ir::MakeAPI(HalideIR::Internal::Stmt, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, tvm::Array<tvm::NodeRef, void>, int, bool)+0x2d85) [0x7f7e649b2c35]
[bt] (3) tvm/build/libtvm.so(+0x22217f) [0x7f7e6485617f]
[bt] (4) tvm/build/libtvm.so(TVMFuncCall+0x5e) [0x7f7e64c292fe]
[bt] (5) tvm/python/tvm/_ffi/_cy2/core.so(+0x1864d) [0x7f7e6037464d]
[bt] (6) tvm/python/tvm/_ffi/_cy2/core.so(+0x18d3b) [0x7f7e60374d3b]
[bt] (7) python(PyEval_EvalFrameEx+0x578f) [0x4c15bf]
[bt] (8) python(PyEval_EvalCodeEx+0x306) [0x4b9ab6]
[bt] (9) python(PyEval_EvalFrameEx+0x58b7) [0x4c16e7]


Error during compile graph
--------------------------
Graph(%input0,
      %input1,
      %input2,
      %input3) {
  %input0, shape=[1,256,5,5]
  %input1, shape=[256,256,2,2]
  %input2, shape=[256,1,1]
  %input3, shape=[256,1,1]
  %2 = conv2d_transpose(%input0, %input1, kernel_size='[2L, 2L]', use_bias='False', dilation='(1, 1)', channels='256', output_padding='(0, 0)', strides='(2, 2)', layout='NCHW', kernel_layout='OIHW', groups='1', padding='(0, 0)'), shape=[1,256,10,10]
  %4 = broadcast_mul(%2, %input2), shape=[1,256,10,10]
  %6 = broadcast_add(%4, %input3), shape=[1,256,10,10]
  %7 = relu(%6), shape=[1,256,10,10]
  ret %7
}
graph_attr_keys = [shape, shape_num_unknown_nodes, dtype, dtype_num_unknown_nodes]

When training SSD model with autoTVM on Mali GPU, the above mistake happened. Can someone help solve this problem?@merrymercy

Can you share the model definition?

Currently we don’t have an optimized conv2d transpose for mali (I can do it later).
Why do you have a conv2d transpose on SSD model?

Thanks! Here is the model. https://github.com/feng98/Model.git

File “/opt/software/tvm/nnvm/python/nnvm/_base.py”, line 75, in check_call
raise NNVMError(py_str(_LIB.NNGetLastError()))
nnvm._base.NNVMError: TVMCall CFunc Error:
Traceback (most recent call last):
File “/opt/software/tvm/python/tvm/_ffi/_ctypes/function.py”, line 54, in cfun
rv = local_pyfunc(*pyargs)
File “/opt/software/tvm/nnvm/python/nnvm/compiler/build_module.py”, line 124, in _build
return tvm.build(funcs, target=target, target_host=target_host)
File “/opt/software/tvm/python/tvm/build_module.py”, line 504, in build
mdev = codegen.build_module(fdevice, str(target_device))
File “/opt/software/tvm/python/tvm/codegen.py”, line 20, in build_module
return _Build(lowered_func, target)
File “/opt/software/tvm/python/tvm/_ffi/function.py”, line 280, in my_api_func
return flocal(*args)
File “/opt/software/tvm/python/tvm/_ffi/_ctypes/function.py”, line 184, in call
ctypes.byref(ret_val), ctypes.byref(ret_tcode)))
File “/opt/software/tvm/python/tvm/_ffi/base.py”, line 66, in check_call
raise TVMError(py_str(_LIB.TVMGetLastError()))
TVMError: [14:56:02] /opt/software/tvm/tvm/src/codegen/codegen.cc:27: Check failed: bf != nullptr Target cuda is not enabled

Do you know this problem ? can you solve it?

This looks like tvm.target.cuda() is being used. You may want to try tvm.target.mali() instead.

I also have this problem when I use conv2d_transpose with the cuda target. Can you help me solve this problem?