masahi commented on a change in pull request #4418: [RUNTIME] Add cudnn conv3d
URL: https://github.com/apache/incubator-tvm/pull/4418#discussion_r351176854
##########
File path: src/runtime/contrib/cudnn/conv_forward.cc
##########
@@ -59,40 +54,125 @@ TVM_REGISTER_GLOBAL("tvm.contrib.cudnn.conv2d.forward")
// Set Data Type
entry_ptr->conv_entry.data_type =
CuDNNDataType::DLTypeToCuDNNType(String2TVMType(conv_dtype));
cudnnDataType_t data_type = CuDNNDataType::DLTypeToCuDNNType(x->dtype);
+ // Dims includes N and C
+ int full_dims = dims + 2;
+
+ std::vector<int> dim_v(full_dims);
+ std::vector<int> tensor_stride_v(full_dims);
+
+ // Note: For 2D tenor, using ND setters causes CUDNN_STATUS_NOT_SUPPORTED
error
+ // in following cudnnGetConvolutionForwardWorkspaceSize() when data type is
fp16, int
+ if (dims == 2) {
// Set Desc
- CUDNN_CALL(cudnnSetConvolution2dDescriptor(entry_ptr->conv_entry.conv_desc,
- pad_h,
- pad_w,
- stride_h,
- stride_w,
- dilation_h,
- dilation_w,
- entry_ptr->conv_entry.mode,
- entry_ptr->conv_entry.data_type));
+ CUDNN_CALL(cudnnSetConvolution2dDescriptor(entry_ptr->conv_entry.conv_desc,
+ pad_v[0],
+ pad_v[1],
+ stride_v[0],
+ stride_v[1],
+ dilation_v[0],
+ dilation_v[1],
+ entry_ptr->conv_entry.mode,
+
entry_ptr->conv_entry.data_type));
+
+ int wn, wc, wh, ww;
+ int xn, xc, xh, xw;
+ int yn, yc, yh, yw;
+ if (entry_ptr->conv_entry.tensor_format == CUDNN_TENSOR_NHWC) {
+ wn = static_cast<int>(w->shape[0]);
+ wc = static_cast<int>(w->shape[3]);
+ wh = static_cast<int>(w->shape[1]);
+ ww = static_cast<int>(w->shape[2]);
+
+ xn = static_cast<int>(x->shape[0]);
+ xc = static_cast<int>(x->shape[3]);
+ xh = static_cast<int>(x->shape[1]);
+ xw = static_cast<int>(x->shape[2]);
+
+ yn = static_cast<int>(y->shape[0]);
+ yc = static_cast<int>(y->shape[3]);
+ yh = static_cast<int>(y->shape[1]);
+ yw = static_cast<int>(y->shape[2]);
+ } else {
+ wn = static_cast<int>(w->shape[0]);
+ wc = static_cast<int>(w->shape[1]);
+ wh = static_cast<int>(w->shape[2]);
+ ww = static_cast<int>(w->shape[3]);
+
+ xn = static_cast<int>(x->shape[0]);
+ xc = static_cast<int>(x->shape[1]);
+ xh = static_cast<int>(x->shape[2]);
+ xw = static_cast<int>(x->shape[3]);
+
+ yn = static_cast<int>(y->shape[0]);
+ yc = static_cast<int>(y->shape[1]);
+ yh = static_cast<int>(y->shape[2]);
+ yw = static_cast<int>(y->shape[3]);
+ }
+
+ // Set Filter
+ CUDNN_CALL(cudnnSetFilter4dDescriptor(entry_ptr->conv_entry.filter_desc,
+ data_type,
+ entry_ptr->conv_entry.tensor_format,
+ wn,
+ wc,
+ wh,
+ ww));
+ // Set Input
+ CUDNN_CALL(cudnnSetTensor4dDescriptor(entry_ptr->conv_entry.input_desc,
+ entry_ptr->conv_entry.tensor_format,
+ data_type,
+ xn,
+ xc,
+ xh,
+ xw));
+ // Set Output
+ CUDNN_CALL(cudnnSetTensor4dDescriptor(entry_ptr->conv_entry.output_desc,
+ entry_ptr->conv_entry.tensor_format,
+ data_type,
+ yn,
+ yc,
+ yh,
+ yw));
+ } else {
+ CUDNN_CALL(cudnnSetConvolutionNdDescriptor(entry_ptr->conv_entry.conv_desc,
+ dims,
+ pad_v,
+ stride_v,
+ dilation_v,
+ entry_ptr->conv_entry.mode,
+
entry_ptr->conv_entry.data_type));
+
// Set Filter
- CUDNN_CALL(cudnnSetFilter4dDescriptor(entry_ptr->conv_entry.filter_desc,
+ for (int i = 0; i < full_dims; i++) {
Review comment:
Please fix the indentation of this block
----------------------------------------------------------------
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
