[Relay] CUDA_ERROR_INVALID_VALUE occurs when testing conv2d_grad

Questions
#1

Hi, I want to test conv2d_grad on GPU, but got some error.

I added this test case in “def test_conv2d_grad()” in /tvm/tests/python/relay/test_op_grad_level2.py
verify_conv2d_grad((1, 736, 17, 17), (128, 736, 1, 1), [1, 1], [0, 0], [1, 1])

when I ran it on GPU, error occurred, this is the error log

Cannot find config for target=cuda, workload=(‘conv2d’, (1, 736, 17, 17, ‘float32’), (128, 736, 1, 1, ‘float32’), (1, 1), (0, 0), (1, 1), ‘NCHW’, ‘float32’). A fallback configuration is used, which may bring great performance regression.
Cannot find config for target=cuda, workload=(‘conv2d_transpose_nchw’, (1, 128, 17, 17, ‘float32’), (128, 736, 1, 1, ‘float32’), (1, 1), (0, 0), ‘float32’). A fallback configuration is used, which may bring great performance regression.
Cannot find config for target=cuda, workload=(‘group_conv2d_nchw’, (1, 736, 17, 17, ‘float32’), (94208, 1, 17, 17, ‘float32’), (1, 1), (0, 0), (1, 1), 736, ‘float32’). A fallback configuration is used, which may bring great performance regression.
Traceback (most recent call last):

File “/tvm/tests/python/relay/test_op_grad_level2.py”, line 132, in
test_conv2d_grad()

File “/tvm/tests/python/relay/test_op_grad_level2.py”, line 125, in test_conv2d_grad
verify_conv2d_grad((1, 736, 17, 17), (128, 736, 1, 1), [1, 1], [0, 0], [1, 1])

File “/tvm/tests/python/relay/test_op_grad_level2.py”, line 113, in verify_conv2d_grad
op_res, (grad_input, grad_weight) = intrp.evaluate(bwd_func)(data, weight)

File “/tvm/python/tvm/relay/backend/interpreter.py”, line 316, in _interp_wrapper
return _intrp(opt_expr)

File “/tvm/python/tvm/_ffi/_ctypes/function.py”, line 210, in call
raise get_last_ffi_error()

tvm._ffi.base.TVMError: Traceback (most recent call last):
[bt] (6) /tvm/build/libtvm.so(TVMFuncCall+0x95) [0x7f5ddc183d8a]
[bt] (5) /tvm/build/libtvm.so(tvm::runtime::PackedFunc::CallPacked(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const+0x30) [0x7f5ddb810ecc]
[bt] (4) /tvm/build/libtvm.so(std::function<void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)>::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const+0x5a) [0x7f5ddb6a82c2]
[bt] (3) /tvm/build/libtvm.so(std::Function_handler<void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), tvm::runtime::detail::PackFuncVoidAddr<4, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocatortvm::runtime::detail::ArgConvertCode > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}>::_M_invoke(std::Any_data const&, tvm::runtime::TVMArgs&&, tvm::runtime::TVMRetValue*&&)+0x58) [0x7f5ddc2051d2]
[bt] (2) /tvm/build/libtvm.so(tvm::runtime::detail::PackFuncVoidAddr<4, tvm::runtime::CUDAWrappedFunc>(tvm::runtime::CUDAWrappedFunc, std::vector<tvm::runtime::detail::ArgConvertCode, std::allocatortvm::runtime::detail::ArgConvertCode > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const+0x23d) [0x7f5ddc2030db]
[bt] (1) /tvm/build/libtvm.so(tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const+0x5ae) [0x7f5ddc20167c]
[bt] (0) /tvm/build/libtvm.so(dmlc::LogMessageFatal::~LogMessageFatal()+0x34) [0x7f5ddb62af2c]
File “/tvm/src/runtime/cuda/cuda_module.cc”, line 215
File “/tvm/src/runtime/module_util.cc”, line 73
TVMError: Check failed: ret == 0 (-1 vs. 0) : CUDALaunch Error: CUDA_ERROR_INVALID_VALUE
grid=(1,94208,1), block=(1,1,1)
// func_name=fused_nn_conv2d_1_kernel0
// CUDA Source
// -----------
extern “C” global void fused_nn_conv2d_1_kernel0( float* restrict placeholder, float* restrict placeholder1, float* restrict compute) {
float compute_local[1];
shared float pad_temp_shared[1];
shared float placeholder_shared[1];
compute_local[0] = 0.000000e+00f;
for (int ry_outer = 0; ry_outer < 17; ++ry_outer) {
for (int rx_outer = 0; rx_outer < 17; ++rx_outer) {
pad_temp_shared[0] = placeholder[((((((int)blockIdx.y) / 128) * 289) + (ry_outer * 17)) + rx_outer)];
placeholder_shared[0] = placeholder1[(((((int)blockIdx.y) * 289) + (ry_outer * 17)) + rx_outer)];
compute_local[0] = (compute_local[0] + (pad_temp_shared[0] * placeholder_shared[0]));
}
}
compute[((int)blockIdx.y)] = compute_local[0];
}
Process finished with exit code 1

It works fine with CPU.
@ vinx13 Could you give me some suggestions?
or Could anyone give me some suggestions? Thanks very much

#2

The default config for group conv is invalid because of too many blocks. You can tune with AutoTVM to get valid configs for operators

#3

Thanks for your reply, it does work.
One problem is that if the parameters of conv changes a little, such as only batch size, we need to tune valid schedule for the new workload again, or error may occur.
Do you have plan to optimize the implementation of conv2d_grad or default schedule for group conv?

Here are some example workloads fromn MXNet InceptionV3, the num_filter parameter of group conv increases to big size in these examples, and the default config doesn’t work for them:

verify_conv2d_grad((10, 80, 73, 73), (192, 80, 3, 3), (1, 1), (0, 0), (1, 1))
verify_conv2d_grad((10, 288, 35, 35), (64, 288, 1, 1), (1, 1), (0, 0), (1, 1))
verify_conv2d_grad((10, 288, 35, 35), (384, 288, 3, 3), (2, 2), (0, 0), (1, 1))
verify_conv2d_grad((10, 768, 17, 17), (192, 768, 1, 1), (1, 1), (0, 0), (1, 1))
verify_conv2d_grad((10, 128, 17, 17), (128, 128, 1, 7), (1, 1), (0, 3), (1, 1))
verify_conv2d_grad((10, 160, 17, 17), (192, 160, 1, 7), (1, 1), (0, 3), (1, 1))
verify_conv2d_grad((10, 1280, 8, 8), (384, 1280, 1, 1), (1, 1), (0, 0), (1, 1))
verify_conv2d_grad((10, 1280, 8, 8), (448, 1280, 1, 1), (1, 1), (0, 0), (1, 1))
verify_conv2d_grad((10, 448, 8, 8), (384, 448, 3, 3), (1, 1), (1, 1), (1, 1))

Hope this will help, thanks.

#4

We can add a fallback config, something like https://github.com/dmlc/tvm/blob/master/topi/python/topi/x86/conv2d.py#L39
This will make sure that default config is valid and you can run using it. But it is difficult to achieve good performance.

#5

OK, I will try that, thank you.

#6

Does auto-tune work for you? It seems that I have similar error

#7

You can refer CUDA C++ Programming Guide . Maximum y- or z-dimension of a grid of thread blocks is 65535.