masahi commented on a change in pull request #4418: [RUNTIME] Add cudnn conv3d
URL: https://github.com/apache/incubator-tvm/pull/4418#discussion_r352031580
##########
File path: src/runtime/contrib/cudnn/conv_forward.cc
##########
@@ -120,137 +180,145 @@ TVM_REGISTER_GLOBAL("tvm.contrib.cudnn.conv2d.forward")
CuDNNDataType::GetConst<0>(entry_ptr->conv_entry.data_type),
entry_ptr->conv_entry.output_desc,
y->data));
-});
+}
-TVM_REGISTER_GLOBAL("tvm.contrib.cudnn.conv2d.output_shape")
-.set_body([](TVMArgs args, TVMRetValue *ret) {
+void OutputShape(
+ int format,
+ const std::vector<int> pad,
+ const std::vector<int> stride,
+ const std::vector<int> dilation,
+ const std::vector<int> x_dim,
+ const std::vector<int> w_dim,
+ void *out_shape,
+ const std::string& data_dtype,
+ const std::string& conv_dtype) {
+ int dims = pad.size();
+
CuDNNThreadEntry* entry_ptr = CuDNNThreadEntry::ThreadLocal();
- int format = args[0];
- int pad_h = args[1];
- int pad_w = args[2];
- int stride_h = args[3];
- int stride_w = args[4];
- int dilation_h = args[5];
- int dilation_w = args[6];
- int x_dim0 = args[7];
- int x_dim1 = args[8];
- int x_dim2 = args[9];
- int x_dim3 = args[10];
- int w_dim0 = args[11];
- int w_dim1 = args[12];
- int w_dim2 = args[13];
- int w_dim3 = args[14];
- void *out_shape = args[15];
- std::string data_dtype = args[16];
- std::string conv_dtype = args[17];
+
// Set Data Type
entry_ptr->conv_entry.data_type =
CuDNNDataType::DLTypeToCuDNNType(String2TVMType(conv_dtype));
cudnnDataType_t data_type =
CuDNNDataType::DLTypeToCuDNNType(String2TVMType(data_dtype));
// Set Format
entry_ptr->conv_entry.tensor_format =
static_cast<cudnnTensorFormat_t>(format);
+ // Dims includes N and C
+ int full_dims = dims + 2;
+
// conv desc
- CUDNN_CALL(cudnnSetConvolution2dDescriptor(entry_ptr->conv_entry.conv_desc,
- pad_h,
- pad_w,
- stride_h,
- stride_w,
- dilation_h,
- dilation_w,
+ CUDNN_CALL(cudnnSetConvolutionNdDescriptor(entry_ptr->conv_entry.conv_desc,
+ dims,
+ pad.data(),
+ stride.data(),
+ dilation.data(),
CUDNN_CROSS_CORRELATION,
entry_ptr->conv_entry.data_type));
- // input desc
- CUDNN_CALL(cudnnSetTensor4dDescriptor(entry_ptr->conv_entry.input_desc,
- entry_ptr->conv_entry.tensor_format,
- data_type,
- x_dim0,
- x_dim1,
- x_dim2,
- x_dim3));
- // filter desc
- CUDNN_CALL(cudnnSetFilter4dDescriptor(entry_ptr->conv_entry.filter_desc,
- data_type,
- entry_ptr->conv_entry.tensor_format,
- w_dim0,
- w_dim1,
- w_dim2,
- w_dim3));
-
-
CUDNN_CALL(cudnnGetConvolution2dForwardOutputDim(entry_ptr->conv_entry.conv_desc,
-
entry_ptr->conv_entry.input_desc,
-
entry_ptr->conv_entry.filter_desc,
-
static_cast<int*>(out_shape),
-
static_cast<int*>(out_shape) + 1,
-
static_cast<int*>(out_shape) + 2,
-
static_cast<int*>(out_shape) + 3));
-});
+ if (dims == 2 && entry_ptr->conv_entry.tensor_format == CUDNN_TENSOR_NHWC) {
+ // Set Input
+ CUDNN_CALL(cudnnSetTensor4dDescriptor(entry_ptr->conv_entry.input_desc,
+ entry_ptr->conv_entry.tensor_format,
+ data_type,
+ x_dim[0],
+ x_dim[3],
+ x_dim[1],
+ x_dim[2]));
+
+ // filter desc
+ CUDNN_CALL(cudnnSetFilter4dDescriptor(entry_ptr->conv_entry.filter_desc,
+ data_type,
+ entry_ptr->conv_entry.tensor_format,
+ w_dim[0],
+ w_dim[3],
+ w_dim[1],
+ w_dim[2]));
+
+
CUDNN_CALL(cudnnGetConvolution2dForwardOutputDim(entry_ptr->conv_entry.conv_desc,
+
entry_ptr->conv_entry.input_desc,
+
entry_ptr->conv_entry.filter_desc,
+
static_cast<int*>(out_shape),
+
static_cast<int*>(out_shape) + 3,
+
static_cast<int*>(out_shape) + 1,
+
static_cast<int*>(out_shape) + 2));
+ } else {
+ // Set Input
+ std::vector<int> tensor_stride(full_dims);
+ GetCudnnStride(full_dims, x_dim, &tensor_stride);
+ CUDNN_CALL(cudnnSetTensorNdDescriptor(entry_ptr->conv_entry.input_desc,
+ data_type,
+ full_dims,
+ x_dim.data(),
+ tensor_stride.data()));
+ // filter desc
+ CUDNN_CALL(cudnnSetFilterNdDescriptor(entry_ptr->conv_entry.filter_desc,
+ data_type,
+ entry_ptr->conv_entry.tensor_format,
+ full_dims,
+ w_dim.data()));
+
+
CUDNN_CALL(cudnnGetConvolutionNdForwardOutputDim(entry_ptr->conv_entry.conv_desc,
+
entry_ptr->conv_entry.input_desc,
+
entry_ptr->conv_entry.filter_desc,
+ full_dims,
+
static_cast<int*>(out_shape)));
+ }
+}
+
+
+void FindAlgo(
+ int format,
+ const std::vector<int> pad,
+ const std::vector<int> stride,
+ const std::vector<int> dilation,
+ const std::vector<int> x_dim,
+ const std::vector<int> w_dim,
+ const std::vector<int> y_dim,
Review comment:
pass by reference
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
[email protected]
With regards,
Apache Git Services
