masahi commented on a change in pull request #4418: [RUNTIME] Add cudnn conv3d
URL: https://github.com/apache/incubator-tvm/pull/4418#discussion_r351186278
##########
File path: src/runtime/contrib/cudnn/conv_forward.cc
##########
@@ -120,137 +200,144 @@ 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,
+ int dims,
+ const int pad_v[],
+ const int stride_v[],
+ const int dilation_v[],
+ const int x_dim_v[],
+ const int w_dim_v[],
+ void *out_shape,
+ const std::string& data_dtype,
+ const std::string& conv_dtype) {
+ // Dims includes N and C
+ int full_dims = dims + 2;
+
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);
+
// 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_v,
+ stride_v,
+ dilation_v,
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) {
Review comment:
why do we need this special case?
----------------------------------------------------------------
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
