Skip to content

[Caffe2] MIOpen bug fixes and performance enhancements#11766

Closed
ashishfarmer wants to merge 11 commits intopytorch:masterfrom
ashishfarmer:af/miopen_fixes
Closed

[Caffe2] MIOpen bug fixes and performance enhancements#11766
ashishfarmer wants to merge 11 commits intopytorch:masterfrom
ashishfarmer:af/miopen_fixes

Conversation

@ashishfarmer
Copy link

This PR contains changes for:

  1. Performance enhancements for group conv using MIOpen
  2. Performance enhancements by removing unnecessary computations while running pooling through MIOpen
  3. Added check for bwdData comptutation while running MIOpen convGradient operator
  4. Fix in MIOpen poolingGradient operator to compute window size for global pooling case
  5. Minor code cleanup in MIOpen spatial batch norm operator

cc: @bddppq @petrex

@petrex
Copy link
Contributor

petrex commented Sep 18, 2018

@ashishfarmer Thanks! Do you have data regarding the performance increase?

@petrex
Copy link
Contributor

petrex commented Sep 18, 2018

@bddppq Is there a way we can run op tests on this particular PR? thx

@bddppq
Copy link
Contributor

bddppq commented Sep 20, 2018

@petrex @ashishfarmer yes you can trigger the tests here https://ci.pytorch.org/jenkins/job/caffe2-builds/job/py2-clang3.8-rocm1.7.1-ubuntu16.04-trigger-test/build?delay=0sec with "GIT_COMMIT" as "origin/pr/YOUR_PR_NUMBER/head" (in this case it's "origin/pr/11766/head").

Triggered here: https://ci.pytorch.org/jenkins/job/caffe2-builds/job/py2-clang3.8-rocm1.7.1-ubuntu16.04-trigger-test/13272/

@bddppq bddppq self-requested a review September 20, 2018 17:42
Copy link
Contributor

@bddppq bddppq left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LG. Could you add a test case that could trigger the issue that you fixed?

beta_(OperatorBase::GetSingleArgument<float>("beta", 0.0)),
do_backward_(
OperatorBase::GetSingleArgument<bool>("do_backward", true)),
OperatorBase::GetSingleArgument<bool>("do_backward", false)),

This comment was marked as off-topic.

This comment was marked as off-topic.

@bddppq
Copy link
Contributor

bddppq commented Sep 20, 2018

Also

vector<int> dims = {N, C, H, W, D};
vector<int> strides = {C * H * W * D, H * W * D, W * D, D, 1};

and
vector<int> dims = {N, C, H, W, D};
vector<int> strides = {C * H * W * D, H * W * D, W * D, D, 1};
look unnecessary to me.

@bddppq
Copy link
Contributor

bddppq commented Sep 20, 2018

And here

MIOPEN_ENFORCE(miopenSet4dTensorDescriptor(
needs to update miopen_input_dims_

@ashishfarmer
Copy link
Author

The case that is fixed in this PR (global pool gradient) is already covered by test_global_pooling in caffe2/python/operator_test/pooling_test.py

@bddppq
Copy link
Contributor

bddppq commented Sep 20, 2018

@ashishfarmer lol then why was it not failing before?

@ashishfarmer
Copy link
Author

It was an intermittent failure before, that we got to the root cause of with this uninitialized variable issue. It is a very specific case (MIOpen engine + AveragePool + NCHW layout) when it will be triggered, and hypothesis would for the most part not encounter it.

@ashishfarmer
Copy link
Author

Also
pytorch/caffe2/operators/hip/spatial_batch_norm_op_miopen.cc

Lines 143 to 144 in 5392b12

vector dims = {N, C, H, W, D};
vector strides = {C * H * W * D, H * W * D, W * D, D, 1};

and
pytorch/caffe2/operators/hip/spatial_batch_norm_op_miopen.cc

Lines 278 to 279 in 5392b12

vector dims = {N, C, H, W, D};
vector strides = {C * H * W * D, H * W * D, W * D, D, 1};
look unnecessary to me.

Removed these vectors

@ashishfarmer
Copy link
Author

And here

pytorch/caffe2/operators/hip/spatial_batch_norm_op_miopen.cc

Line 280 in 5392b12

MIOPEN_ENFORCE(miopenSet4dTensorDescriptor(
needs to update miopen_input_dims_

Good catch! Thank you! Fixed this

@ashishfarmer
Copy link
Author

@bddppq - Added the fixes as per your review

Copy link
Contributor

@facebook-github-bot facebook-github-bot left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

facebook-github-bot pushed a commit that referenced this pull request Sep 20, 2018
Summary:
This PR contains changes for:
1. Performance enhancements for group conv using MIOpen
2. Performance enhancements by removing unnecessary computations while running pooling through MIOpen
3. Added check for bwdData comptutation while running MIOpen convGradient operator
4. Fix in MIOpen poolingGradient operator to compute window size for global pooling case
5. Minor code cleanup in MIOpen spatial batch norm operator

Differential Revision: D9979050

Pulled By: bddppq

fbshipit-source-id: fabc7a44a2f9ca0307d99564d1ce8fe1de9a6fbb
@bddppq
Copy link
Contributor

bddppq commented Sep 21, 2018

Merged to master in c7751f4

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants