Refactor cudnn convolution by zasdfgbnm · Pull Request #49109 · pytorch/pytorch

zasdfgbnm · 2020-12-09T20:50:27Z

cuDNN v7 API has been deprecated, so we need to migrate to cuDNN v8 API. The v8 API does not exist on cuDNN 7, so there will be a long time both API should exist.

This is step 0 of adding cuDNN v8 API. There is no real code change in this PR. It just copy-pastes existing code. The original Conv.cpp is split into ConvPlaceholders.cpp, ConvShared.cpp, ConvShared.h, Conv_v7.cpp, Conv_v8.cpp. Currently Conv_v8.cpp is empty, and will be filled in the future.

The ConvPlaceholders.cpp contains placeholder implementation of cudnn convolution when cudnn is not enabled. These operators only raise errors and do no real computation. This file also contains deprecated operators. These operators are implemented using current operators.

The ConvShared.cpp and ConvShared.h contains code that will be shared by the v7 and v8 API, these include the definition of struct ConvolutionParams and ConvolutionArgs. As well as ATen exposed API like cudnn_convolution and intermediate cudnn_convolution_forward. These exposed functions will call raw API like raw_cudnn_convolution_forward_out in Conv_v7.cpp or Conv_v8.cpp for the real implementation.

The Conv_v7.cpp, Conv_v8.cpp contains the implementation of raw APIs, and are different for v7 and v8.

dr-ci · 2020-12-09T21:09:11Z

💊 CI failures summary and remediations

As of commit e4be821 (more details on the Dr. CI page):

1/1 failures possibly* introduced in this PR
- 1/1 non-CircleCI failure(s)

Extra GitHub checks: 1 failed

Failed: GitHub Actions - clang-tidy

This comment was automatically generated by Dr. CI (expand for details).

Follow this link to opt-out of these comments for your Pull Requests.

Please report bugs/suggestions on the GitHub issue tracker or post in the (internal) Dr. CI Users group.

See how this bot performed.

This comment has been revised 11 times.

codecov · 2020-12-10T01:57:08Z

Codecov Report

Merging #49109 (e4be821) into master (3f9ff48) will decrease coverage by 0.00%.
The diff coverage is n/a.

@@            Coverage Diff             @@
##           master   #49109      +/-   ##
==========================================
- Coverage   80.73%   80.73%   -0.01%     
==========================================
  Files        1871     1871              
  Lines      201759   201759              
==========================================
- Hits       162888   162885       -3     
- Misses      38871    38874       +3

facebook-github-bot

@ezyang has imported this pull request. If you are a Facebook employee, you can view this diff on Phabricator.

ngimel · 2020-12-10T17:03:15Z

+  params->dataType = dataType;
+  // ASSERT(weight.dim() == input.dim())
+  for (int i = 0; i != input.dim(); ++i) {
+    params->input_size[i] = (int) input.size(i);


Can you please do unchecked accesses here? input.sizes()[i]

will fix in later PR

IIRC, these were all from the old code, right?

Yes, they are, but that's not a reason not to fix them.

ngimel · 2020-12-10T17:12:17Z

+
+  // Input
+  checkDimRange(c, input, 3, 6 /* exclusive */);
+  checkSize(c, input, input_channels_dim, weight->size(1) * groups);


unchecked access for size(1)

ngimel · 2020-12-10T17:13:14Z

+
+  auto layout = cudnn_conv_use_channels_last(*input, *weight) ?
+      at::MemoryFormat::ChannelsLast : at::MemoryFormat::Contiguous;
+  auto output_t = at::empty(


you can do at::native::empty_cuda here to avoid dispatch

ngimel · 2020-12-10T17:24:18Z

+
+  Tensor grad_input, grad_weight;
+  if (output_mask[0]) {
+    grad_input = at::cudnn_convolution_transpose_backward_input(grad_output, weight, padding, stride, dilation, groups, benchmark, deterministic, allow_tf32);


Maybe you could call cudnn_convolution_forward directly here, and get rid of cudnn_convolution and cudnn_convolution_transpose_backward_input? Do they provide any value?

I think the value is a better error message, i.e. correctly calling it grad_output instead of input

facebook-github-bot · 2020-12-10T19:18:30Z

@ezyang merged this pull request in 45473ff.

Summary: - Resolves ngimel's review comments in #49109 - Move `ConvolutionArgs` from `ConvShared.h` to `Conv_v7.cpp`, because cuDNN v8 uses different descriptors therefore will not share the same `ConvolutionArgs`. - Refactor the `ConvolutionParams` (the hash key for benchmark): - Remove `input_stride` - Add `input_dim` - Add `memory_format` - Make `repro_from_args` to take `ConvolutionParams` instead of `ConvolutionArgs` as arguments so that it can be shared for v7 and v8 - Rename some `layout` to `memory_format`. `layout` should be sparse/strided and `memory_format` should be contiguous/channels_last. They are different things. Pull Request resolved: #50827 Reviewed By: bdhirsh Differential Revision: D26048274 Pulled By: ezyang fbshipit-source-id: f71aa02d90ffa581c17ab05b171759904b311517

Summary: cuDNN v7 API has been deprecated, so we need to migrate to cuDNN v8 API. The v8 API does not exist on cuDNN 7, so there will be a long time both API should exist. This is step 0 of adding cuDNN v8 API. There is no real code change in this PR. It just copy-pastes existing code. The original `Conv.cpp` is split into `ConvPlaceholders.cpp`, `ConvShared.cpp`, `ConvShared.h`, `Conv_v7.cpp`, `Conv_v8.cpp`. Currently `Conv_v8.cpp` is empty, and will be filled in the future. The `ConvPlaceholders.cpp` contains placeholder implementation of cudnn convolution when cudnn is not enabled. These operators only raise errors and do no real computation. This file also contains deprecated operators. These operators are implemented using current operators. The `ConvShared.cpp` and `ConvShared.h` contains code that will be shared by the v7 and v8 API, these include the definition of struct `ConvolutionParams` and `ConvolutionArgs`. As well as ATen exposed API like `cudnn_convolution` and intermediate `cudnn_convolution_forward`. These exposed functions will call raw API like `raw_cudnn_convolution_forward_out` in `Conv_v7.cpp` or `Conv_v8.cpp` for the real implementation. The `Conv_v7.cpp`, `Conv_v8.cpp` contains the implementation of raw APIs, and are different for v7 and v8. Pull Request resolved: pytorch#49109 Reviewed By: H-Huang Differential Revision: D25463783 Pulled By: ezyang fbshipit-source-id: 1c80de8e5d94d97a61e45687f6193e8ff5481e3e

Summary: - Resolves ngimel's review comments in pytorch#49109 - Move `ConvolutionArgs` from `ConvShared.h` to `Conv_v7.cpp`, because cuDNN v8 uses different descriptors therefore will not share the same `ConvolutionArgs`. - Refactor the `ConvolutionParams` (the hash key for benchmark): - Remove `input_stride` - Add `input_dim` - Add `memory_format` - Make `repro_from_args` to take `ConvolutionParams` instead of `ConvolutionArgs` as arguments so that it can be shared for v7 and v8 - Rename some `layout` to `memory_format`. `layout` should be sparse/strided and `memory_format` should be contiguous/channels_last. They are different things. Pull Request resolved: pytorch#50827 Reviewed By: bdhirsh Differential Revision: D26048274 Pulled By: ezyang fbshipit-source-id: f71aa02d90ffa581c17ab05b171759904b311517

zasdfgbnm added 2 commits December 9, 2020 11:38

Refactor cudnn convolution

78737b4

fix

11895f9

facebook-github-bot added the cla signed label Dec 9, 2020

pytorchbot added the open source label Dec 9, 2020

Update ConvShared.h

188f5f4

zasdfgbnm requested review from ezyang, ngimel and xwang233 December 9, 2020 21:14

fix

e4be821

facebook-github-bot reviewed Dec 10, 2020

View reviewed changes

ezyang approved these changes Dec 10, 2020

View reviewed changes

ngimel reviewed Dec 10, 2020

View reviewed changes

facebook-github-bot closed this in 45473ff Dec 10, 2020

zasdfgbnm deleted the cudnn7-params-refactor branch December 10, 2020 19:09

facebook-github-bot added the Merged label Dec 10, 2020

zasdfgbnm mentioned this pull request Jan 20, 2021

More about cudnn refactor #50827

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Refactor cudnn convolution#49109

Refactor cudnn convolution#49109
zasdfgbnm wants to merge 4 commits intomasterfrom
cudnn7-params-refactor

zasdfgbnm commented Dec 9, 2020

Uh oh!

dr-ci Bot commented Dec 9, 2020 •

edited

Loading

Uh oh!

codecov Bot commented Dec 10, 2020

Uh oh!

facebook-github-bot left a comment

Uh oh!

ngimel Dec 10, 2020

Uh oh!

zasdfgbnm Dec 10, 2020

Uh oh!

ezyang Dec 10, 2020

Uh oh!

ngimel Dec 11, 2020

Uh oh!

ngimel Dec 10, 2020

Uh oh!

ngimel Dec 10, 2020

Uh oh!

ngimel Dec 10, 2020

Uh oh!

zasdfgbnm Dec 10, 2020

Uh oh!

facebook-github-bot commented Dec 10, 2020

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

zasdfgbnm commented Dec 9, 2020

Uh oh!

dr-ci Bot commented Dec 9, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

💊 CI failures summary and remediations

Extra GitHub checks: 1 failed

Uh oh!

codecov Bot commented Dec 10, 2020

Codecov Report

Uh oh!

facebook-github-bot left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

facebook-github-bot commented Dec 10, 2020

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

dr-ci Bot commented Dec 9, 2020 •

edited

Loading