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)