masahi commented on a change in pull request #4418: [RUNTIME] Add cudnn conv3d
URL: https://github.com/apache/incubator-tvm/pull/4418#discussion_r351180666
##########
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]);
+ }
+
Review comment:
I think you can clean up this. You can setup batch, channel, h, w indices
based on layout, and then access shapes with these indices.
----------------------------------------------------------------
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
