[CUDA]Got Error: CUDA ERROR LAUNCH OUT OF RESOURCES

blacklong617 · October 8, 2019, 10:57am

I am preparing to do float16 inference with tvm, I load a resnet-18 model from onnx,the code is

onnx_model = onnx.load('models/resnet18_half.onnx')
input_name = '0'
x = np.random.uniform(-1, 1, size=(1,3,256,340)).astype("float16")

shape_dict = {input_name: x.shape}
sym, params = relay.frontend.from_onnx(onnx_model, shape_dict,dtype="float16")
opt_level = 3
target = tvm.target.cuda()
network = 'resnet18-rtx2080-fp16'
log_file = "%s.log" % network
from tvm import autotvm
with autotvm.apply_history_best(log_file):
    with relay.build_config(opt_level=opt_level):
        graph, lib, params = relay.build_module.build(
            sym, target, params=params)

    print(sym.astext(show_meta_data=True))
    # create random input
    ctx = tvm.gpu()
    module = graph_runtime.create(graph, lib, ctx)
    # set input and parameters

    module.set_input(**params)
    print("start loop!")
    # evaluate
    module.set_input("0", x)
    print("Evaluate inference time cost...")
    module.run()

I got a error like:

  Traceback (most recent call last):

  File "onnx_resnet18_fp16.py", line 56, in <module>
    module.run()

  File "/home/xxx/project/tvm-debug/python/tvm/contrib/graph_runtime.py", line 168, in run
    self._run()

  File "/home/xxx/project/tvm-debug/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] (3) /home/xxx/project/tvm-debug/build/libtvm.so(TVMFuncCall+0x61) [0x7f86ff18a821]
  [bt] (2) /home/xxx/project/tvm-debug/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::allocator<tvm::runtime::detail::ArgConvertCode> > const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}>::_M_invoke(std::_Any_data const&, tvm::runtime::TVMArgs&&, tvm::runtime::TVMRetValue*&&)+0xbc) [0x7f86ff1f0f5c]
  [bt] (1) /home/xxx/project/tvm-debug/build/libtvm.so(tvm::runtime::CUDAWrappedFunc::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*, void**) const+0x662) [0x7f86ff1f09e2]
  [bt] (0) /home/xxx/project/tvm-debug/build/libtvm.so(dmlc::LogMessageFatal::~LogMessageFatal()+0x32) [0x7f86fe9ea1f2]
  File "/home/xxx/project/tvm-debug/src/runtime/cuda/cuda_module.cc", line 215
  File "/home/xxx/project/tvm-debug/src/runtime/module_util.cc", line 73
TVMError: Check failed: ret == 0 (-1 vs. 0) : CUDALaunch Error: CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES
 grid=(1,1,1),  block=(1024,1,1)
// func_name=fused_nn_avg_pool2d_kernel0
// CUDA Source
// -----------
#include <cuda_fp16.h>
__device__ half max(const half a, const half b)
{
  return __hlt(a, b) ? b : a;
}
__device__ half min(const half a, const half b)
{
  return __hlt(__half(b), __half(a)) ? b : a;
}
__device__ half operator + (const volatile __half &a,  const volatile __half &b)
{
  return __hadd(::__half(a), ::__half(b));
}

extern "C" __global__ void fused_nn_conv2d_expand_dims_multiply_negative_multiply_add_expand_dims_add_2_kernel0( half* __restrict__ placeholder,  half* __restrict__ placeholder1,  half* __restrict__ T_add,  half* __restrict__ placeholder2,  half* __restrict__ placeholder3,  half* __restrict__ placeholder4) {
   half compute[4];
  __shared__ half pad_temp_shared[1008];
  __shared__ half placeholder_shared[1024];
  #pragma unroll
  for (int yy_init = 0; yy_init < 2; ++yy_init) {
    compute[yy_init] = __float2half_rn(0.000000e+00f);
    compute[(yy_init + 2)] = __float2half_rn(0.000000e+00f);
  }
  #pragma unroll
  for (int rc_outer = 0; rc_outer < 16; ++rc_outer) {
    __syncthreads();
    #pragma unroll
    for (int ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner = 0; ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner < 3; ++ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) {
      if ((((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) < 1008) {
        if (((((int)threadIdx.x) * 3) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) < 32) {
          pad_temp_shared[(((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner)] = placeholder[(((((rc_outer * 5632) + (((((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) / 63) * 352)) + (((int)blockIdx.y) * 88)) + ((((((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) % 63) / 21) * 22)) + ((((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner) % 21))];
        }
      }
    }
    #pragma unroll
    for (int ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1 = 0; ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1 < 3; ++ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) {
      if (((((int)threadIdx.z) * 2) + (((((int)threadIdx.x) * 3) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) / 16)) < 64) {
        if ((((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) < 1024) {
          if (((((int)threadIdx.x) * 3) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) < 32) {
            placeholder_shared[(((((int)threadIdx.z) * 32) + (((int)threadIdx.x) * 3)) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1)] = placeholder1[(((((((int)blockIdx.z) * 16384) + (((int)threadIdx.z) * 512)) + ((((((int)threadIdx.x) * 3) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) / 16) * 256)) + (rc_outer * 16)) + (((((int)threadIdx.x) * 3) + ax0_ax1_fused_ax2_fused_ax3_fused_inner_inner_inner1) % 16))];
          }
        }
      }
    }
    __syncthreads();
    #pragma unroll
    for (int rc_inner = 0; rc_inner < 16; ++rc_inner) {
      #pragma unroll
      for (int yy = 0; yy < 2; ++yy) {
        compute[yy] = (compute[yy] + (pad_temp_shared[(((rc_inner * 63) + (yy * 42)) + (((int)threadIdx.x) * 2))] * placeholder_shared[((((int)threadIdx.z) * 16) + rc_inner)]));
        compute[(yy + 2)] = (compute[(yy + 2)] + (pad_temp_shared[(((rc_inner * 63) + (yy * 42)) + (((int)threadIdx.x) * 2))] * placeholder_shared[(((((int)threadIdx.z) * 16) + rc_inner) + 512)]));
      }
    }
  }
  #pragma unroll
  for (int ax2_inner_inner_inner = 0; ax2_inner_inner_inner < 2; ++ax2_inner_inner_inner) {
    T_add[(((((((int)blockIdx.z) * 5632) + (((int)threadIdx.z) * 88)) + (((int)blockIdx.y) * 22)) + (ax2_inner_inner_inner * 11)) + ((int)threadIdx.x))] = ((compute[ax2_inner_inner_inner] * placeholder2[((((int)blockIdx.z) * 64) + ((int)threadIdx.z))]) + (((__float2half_rn(0.000000e+00f) - placeholder3[((((int)blockIdx.z) * 64) + ((int)threadIdx.z))]) * placeholder2[((((int)blockIdx.z) * 64) + ((int)threadIdx.z))]) + placeholder4[((((int)blockIdx.z) * 64) + ((int)threadIdx.z))]));
    T_add[((((((((int)blockIdx.z) * 5632) + (((int)threadIdx.z) * 88)) + (((int)blockIdx.y) * 22)) + (ax2_inner_inner_inner * 11)) + ((int)threadIdx.x)) + 2816)] = ((compute[(ax2_inner_inner_inner + 2)] * placeholder2[(((((int)blockIdx.z) * 64) + ((int)threadIdx.z)) + 32)]) + (((__float2half_rn(0.000000e+00f) - placeholder3[(((((int)blockIdx.z) * 64) + ((int)threadIdx.z)) + 32)]) * placeholder2[(((((int)blockIdx.z) * 64) + ((int)threadIdx.z)) + 32)]) + placeholder4[(((((int)blockIdx.z) * 64) + ((int)threadIdx.z)) + 32)]));
  }
}

extern "C" __global__ void fused_nn_conv2d_expand_dims_multiply_negative_multiply_add_expand_dims_add_nn_re_16196308750733684494__3_kernel0( half* __restrict__ placeholder,  half* __restrict__ placeholder1,  half* __restrict__ T_relu,  half* __restrict__ placeholder2,  half* __restrict__ placeholder3,  half* __restrict__ placeholder4) {
   half compute[32];
  __shared__ half pad_temp_shared[3240];
  __shared__ half placeholder_shared[2304];
  compute[0] = __float2half_rn(0.000000e+00f);
  compute[16] = __float2half_rn(0.000000e+00f);
  compute[8] = __float2half_rn(0.000000e+00f);
  compute[24] = __float2half_rn(0.000000e+00f);
  compute[1] = __float2half_rn(0.000000e+00f);
  compute[17] = __float2half_rn(0.000000e+00f);
  compute[9] = __float2half_rn(0.000000e+00f);
  compute[25] = __float2half_rn(0.000000e+00f);
  compute[2] = __float2half_rn(0.000000e+00f);
  compute[18] = __float2half_rn(0.000000e+00f);
  compute[10] = __float2half_rn(0.000000e+00f);
  compute[26] = __float2half_rn(0.000000e+00f);
  compute[3] = __float2half_rn(0.000000e+00f);
  compute[19] = __float2half_rn(0.000000e+00f);
  compute[11] = __float2half_rn(0.000000e+00f);
  compute[27] = __float2half_rn(0.000000e+00f);
  compute[4] = __float2half_rn(0.000000e+00f);
  compute[20] = __float2half_rn(0.000000e+00f);
  compute[12] = __float2half_rn(0.000000e+00f);
  compute[28] = __float2half_rn(0.000000e+00f);
  compute[5] = __float2half_rn(0.000000e+00f);
  compute[21] = __float2half_rn(0.000000e+00f);
  compute[13] = __float2half_rn(0.000000e+00f);
  compute[29] = __float2half_rn(0.000000e+00f);
  compute[6] = __float2half_rn(0.000000e+00f);
  compute[22] = __float2half_rn(0.000000e+00f);
  compute[14] = __float2half_rn(0.000000e+00f);
  compute[30] = __float2half_rn(0.000000e+00f);
  compute[7] = __float2half_rn(0.000000e+00f);
  compute[23] = __float2half_rn(0.000000e+00f);
  compute[15] = __float2half_rn(0.000000e+00f);
  compute[31] = __float2half_rn(0.000000e+00f);
  for (int rc_outer = 0; rc_outer < 16; ++rc_outer) {
    __syncthreads();
    pad_temp_shared[(((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37))] = ((((1 <= ((((int)blockIdx.y) * 8) + ((((int)threadIdx.x) * 37) / 45))) && (1 <= ((((int)threadIdx.x) * 37) % 45))) && (((((int)threadIdx.x) * 37) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + (((((int)threadIdx.x) * 37) / 45) * 43)) + ((((int)threadIdx.x) * 37) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 1)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 1) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 1) % 45))) && ((((((int)threadIdx.x) * 37) + 1) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 1) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 1) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 2)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 2) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 2) % 45))) && ((((((int)threadIdx.x) * 37) + 2) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 2) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 2) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 3)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 3) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 3) % 45))) && ((((((int)threadIdx.x) * 37) + 3) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 3) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 3) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 4)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 4) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 4) % 45))) && ((((((int)threadIdx.x) * 37) + 4) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 4) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 4) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 5)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 5) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 5) % 45))) && ((((((int)threadIdx.x) * 37) + 5) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 5) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 5) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 6)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 6) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 6) % 45))) && ((((((int)threadIdx.x) * 37) + 6) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 6) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 6) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 7)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 7) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 7) % 45))) && ((((((int)threadIdx.x) * 37) + 7) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 7) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 7) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 8)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 8) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 8) % 45))) && ((((((int)threadIdx.x) * 37) + 8) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 8) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 8) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 9)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 9) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 9) % 45))) && ((((((int)threadIdx.x) * 37) + 9) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 9) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 9) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 10)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 10) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 10) % 45))) && ((((((int)threadIdx.x) * 37) + 10) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 10) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 10) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 11)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 11) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 11) % 45))) && ((((((int)threadIdx.x) * 37) + 11) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 11) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 11) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 12)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 12) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 12) % 45))) && ((((((int)threadIdx.x) * 37) + 12) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 12) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 12) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 13)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 13) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 13) % 45))) && ((((((int)threadIdx.x) * 37) + 13) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 13) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 13) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 14)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 14) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 14) % 45))) && ((((((int)threadIdx.x) * 37) + 14) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 14) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 14) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 15)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 15) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 15) % 45))) && ((((((int)threadIdx.x) * 37) + 15) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 15) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 15) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 16)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 16) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 16) % 45))) && ((((((int)threadIdx.x) * 37) + 16) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 16) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 16) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 17)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 17) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 17) % 45))) && ((((((int)threadIdx.x) * 37) + 17) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 17) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 17) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 18)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 18) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 18) % 45))) && ((((((int)threadIdx.x) * 37) + 18) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 18) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 18) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 19)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 19) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 19) % 45))) && ((((((int)threadIdx.x) * 37) + 19) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 19) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 19) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 20)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 20) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 20) % 45))) && ((((((int)threadIdx.x) * 37) + 20) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 20) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 20) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 21)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 21) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 21) % 45))) && ((((((int)threadIdx.x) * 37) + 21) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 21) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 21) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 22)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 22) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 22) % 45))) && ((((((int)threadIdx.x) * 37) + 22) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 22) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 22) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 23)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 23) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 23) % 45))) && ((((((int)threadIdx.x) * 37) + 23) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 23) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 23) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 24)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 24) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 24) % 45))) && ((((((int)threadIdx.x) * 37) + 24) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 24) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 24) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 25)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 25) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 25) % 45))) && ((((((int)threadIdx.x) * 37) + 25) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 25) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 25) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 26)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 26) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 26) % 45))) && ((((((int)threadIdx.x) * 37) + 26) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 26) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 26) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 27)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 27) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 27) % 45))) && ((((((int)threadIdx.x) * 37) + 27) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 27) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 27) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 28)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 28) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 28) % 45))) && ((((((int)threadIdx.x) * 37) + 28) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 28) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 28) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 29)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 29) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 29) % 45))) && ((((((int)threadIdx.x) * 37) + 29) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 29) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 29) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 30)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 30) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 30) % 45))) && ((((((int)threadIdx.x) * 37) + 30) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 30) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 30) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 31)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 31) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 31) % 45))) && ((((((int)threadIdx.x) * 37) + 31) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 31) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 31) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 32)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 32) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 32) % 45))) && ((((((int)threadIdx.x) * 37) + 32) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 32) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 32) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 33)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 33) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 33) % 45))) && ((((((int)threadIdx.x) * 37) + 33) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 33) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 33) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 34)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 34) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 34) % 45))) && ((((((int)threadIdx.x) * 37) + 34) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 34) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 34) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
    if ((((((int)threadIdx.z) * 2) + (((((int)threadIdx.x) * 37) + 35) / 405)) + ((int)threadIdx.y)) < 8) {
      if ((((((int)threadIdx.z) * 18) + (((int)threadIdx.y) * 9)) + (((((int)threadIdx.x) * 37) + 35) / 45)) < 72) {
        if ((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) < 3205) {
          if (((((int)threadIdx.y) * 405) + (((int)threadIdx.x) * 37)) < 775) {
            if (((int)threadIdx.x) < 10) {
              pad_temp_shared[((((((int)threadIdx.z) * 810) + (((int)threadIdx.y) * 405)) + (((int)threadIdx.x) * 37)) + 35)] = ((((1 <= ((((int)blockIdx.y) * 8) + (((((int)threadIdx.x) * 37) + 35) / 45))) && (1 <= (((((int)threadIdx.x) * 37) + 35) % 45))) && ((((((int)threadIdx.x) * 37) + 35) % 45) < 44)) ? placeholder[(((((((rc_outer * 11008) + (((int)threadIdx.z) * 2752)) + (((int)threadIdx.y) * 1376)) + (((int)blockIdx.y) * 344)) + ((((((int)threadIdx.x) * 37) + 35) / 45) * 43)) + (((((int)threadIdx.x) * 37) + 35) % 45)) - 44)] : __float2half_rn(0.000000e+00f));
            }
          }
        }
      }
    }...

My cuda version is 10.0
My gpu is RTX2080
Who can help me see what is the cause of the error?
Interestingly, I can run the code well on my gtx1060 gpu. Even though the gtx1060 FP16 computing power is too low.(68.36 GFLOPS)

blacklong617 · October 8, 2019, 11:26am

nvidia 2080 super device info:

 ./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "Graphics Device"
  CUDA Driver Version / Runtime Version          10.2 / 10.0
  CUDA Capability Major/Minor version number:    7.5
  Total amount of global memory:                 7981 MBytes (8368685056 bytes)
  (48) Multiprocessors, ( 64) CUDA Cores/MP:     3072 CUDA Cores
  GPU Max Clock rate:                            1815 MHz (1.81 GHz)
  Memory Clock rate:                             7751 Mhz
  Memory Bus Width:                              256-bit
  L2 Cache Size:                                 4194304 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  1024
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 3 copy engine(s)
  Run time limit on kernels:                     No
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 129 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 10.2, CUDA Runtime Version = 10.0, NumDevs = 1
Result = PASS

nvidia 1060 GPU info:

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 1060 6GB"
  CUDA Driver Version / Runtime Version          10.0 / 10.0
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 6075 MBytes (6370295808 bytes)
  (10) Multiprocessors, (128) CUDA Cores/MP:     1280 CUDA Cores
  GPU Max Clock rate:                            1709 MHz (1.71 GHz)
  Memory Clock rate:                             4004 Mhz
  Memory Bus Width:                              192-bit
  L2 Cache Size:                                 1572864 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 10.0, CUDA Runtime Version = 10.0, NumDevs = 1
Result = PASS

comaniac · October 8, 2019, 4:45pm

You may encounter this error when the schedule costs too much device memory, so this is more likely to happen on small GPUs. Is your log file generated by AutoTVM?

blacklong617 · October 9, 2019, 1:27am

The error happened when I run “module.run()”.
And it works well on GTX1060 and T4 but got error on RTX2080.
I have read the error type :

CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701
This indicates that a launch did not occur because it did not have appropriate resources. This error
usually indicates that the user has attempted to pass too many arguments to the device kernel, or the
kernel launch specifies too many threads for the kernel's register count. Passing arguments of the 
wrong size (i.e. a 64-bit pointer when a 32-bit int is expected) is equivalent to passing too many 
arguments and can also result in this error.

I guess this error is related to the FP16 I added about cuda. :

TVMError: Check failed: ret == 0 (-1 vs. 0) : CUDALaunch Error: CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES
 grid=(1,1,1),  block=(1024,1,1)
// func_name=fused_nn_avg_pool2d_kernel0
// CUDA Source
// -----------
#include <cuda_fp16.h>
__device__ half max(const half a, const half b)
{
  return __hgt(__half(a), __half(b)) ? a : b;
}
__device__ half min(const half a, const half b)
{
  return __hgt(__half(a), __half(b)) ? a : b;
}
__device__ half operator+(const volatile __half &a,  const volatile __half &b)
{
  return __hadd(a, b);
}

But why can works well on GTX1060 ,I do not know.
I saw that the difference between the two of them(GTX1060 RTX2080) is obvious.:
Maximum number of threads per multiprocessor: 2048 (GTX1060)
Maximum number of threads per multiprocessor: 1024 (RTX2080)
Who can answer for me what is causing such a result ?
Thanks!

blacklong617 · October 22, 2019, 2:45am

The onnx model I used resnet-18-fp16.onnx
The complete output log error.log

tiandiao123 · October 20, 2020, 7:35pm

In some sense, I decreased batch size, it usually worked out. Did you find some solution to solve this problem? I mean let module.run() without getting out of cuda memory error?

peaceblues · April 1, 2021, 10:29am

Is your issue fixed? I was trapped in the same issue, opt_level set to 0/1, it runs well, but opt_level >=2, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES show up when “module.run()”

with tvm.transform.PassContext(opt_level=3):

lib = relay.build(mod, target=target, params=params)

my gpu is Tesla V100 smx2.

any lessson could learned from you?

@comaniac @tiandiao123 @apivovarov @merrymercy @eqy