You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mxnet.apache.org by GitBox <gi...@apache.org> on 2022/07/22 03:34:51 UTC
[GitHub] [incubator-mxnet] DickJC123 opened a new pull request, #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
DickJC123 opened a new pull request, #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104
## Description ##
PR #20635, which began using the cuDNN v8 backend API for Convolution ops, includes the following line to avoid `test_gluon_gpu.py::test_group_conv2d_16c` failures that began occurring coincident with the PR:
```
@with_environment('MXNET_CUDNN_DISABLED_CONV_FWD_ENGINES', '5') # eng:5 causes test failure on M60
```
This PR will remove that line by providing a different implementation of the "convolution plan cache" introduced with PR #20635 that is compatible with convolution engine 5. The steps of this PR will be:
1. Reenable convolution engine 5, and demonstrate a return of `test_group_conv2d` failures, then
2. Add the upgrade to the convolution plan cache so that `test_group_conv2d` passes even with engine 5 use
Further detail:
The cuDNN v8 backend allows one to bypass a lot of CPU processing that might precede kernel launch by first building up and finalizing a convolution execution plan. The plan is then executed efficiently by the call `cudnnBackendExecute(cudnn_handle, plan, ...)`. PR #20635 introduced a cache of plans so that autotuning does not need to be repeated for identically-parameterized convolutions, and so that such convolutions are handled by the same engine even if they exist multiple times in a model or are handled by different GPU workers.
The issue that was discovered for convolution engine 5 is that it caches a cuDNN handle provided during the plan's construction, and does not consider the handle passed as an argument of cudnnBackendExecute(). The result is that the engine's kernels are launched into the stream of the cached handle, and this would be the incorrect stream if the GPU worker launching the plan is different from the one that created the plan. Without the proper stream synchronization, incorrect results may follow.
The contribution of this PR is to effectively include a GPU worker's cudnn handle as part of the key used in the cache lookup. One aspect of the fix though is that if there's a cache miss, an entry made by a different worker can be 'cloned' with the proper handle without repeating the autotuning.
## Checklist ##
### Essentials ###
- [X] PR's title starts with a category (e.g. [BUGFIX], [MODEL], [TUTORIAL], [FEATURE], [DOC], etc)
- [ ] Changes are complete (i.e. I finished coding on this PR)
- [X] All changes have test coverage
- [X] Code is well-documented
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] DickJC123 commented on a diff in pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
DickJC123 commented on code in PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#discussion_r936078159
##########
src/operator/cudnn_ops.h:
##########
@@ -121,25 +128,39 @@ void MaybeLogSelectedPlan(const cudnn_cxx::Descriptor& plan);
// Op::Param - a type, collecting all data, required to create cuDNN descriptor(s), but not needed
// for execution.
// Op::MakeKey() - a static function, which maps its arguments to a tuple - a key in the op cache.
-// Op::Make() - a static function, creating the necessary cuDNN descriptor.
-// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor and
+// Op::Make() - a static function, creating all necessary cuDNN descriptors.
+// Op::Clone() - a static function, creating a copy of the op's descriptors with a new cudNN handle.
+// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor(s) and
// the passed arguments.
template <typename Op, typename... Args>
bool Exec(const OpContext& ctx, const typename Op::Param& param, Args&&... args) {
auto key = std::tuple_cat(std::make_tuple(ctx.run_ctx.ctx.dev_id),
Op::MakeKey(param, std::forward<Args>(args)...));
- static std::unordered_map<decltype(key), cudnn_cxx::Descriptor> op_map;
static std::mutex mx;
std::unique_lock<std::mutex> lk(mx);
- auto it = op_map.find(key);
- if (it == op_map.end()) {
- auto op = Op::Make(ctx, param, std::forward<Args>(args)...);
- it = op_map.emplace(key, std::move(op)).first;
- }
+ static std::unordered_multimap<decltype(key), const cudnn_cxx::Descriptor> op_map;
+ auto match_it = [&]() {
+ // Some cuDNN Op implementations require that the thread's cuDNN handle
+ // (used in cudnnBackendExecute()) matches the one used in making the plan.
+ const bool ignore_handles = false;
Review Comment:
Right, but the plan is eventually to change `ignore_handles` to be a function of CUDNN_VERSION, once the proper version is known.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] DickJC123 commented on a diff in pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
DickJC123 commented on code in PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#discussion_r936078159
##########
src/operator/cudnn_ops.h:
##########
@@ -121,25 +128,39 @@ void MaybeLogSelectedPlan(const cudnn_cxx::Descriptor& plan);
// Op::Param - a type, collecting all data, required to create cuDNN descriptor(s), but not needed
// for execution.
// Op::MakeKey() - a static function, which maps its arguments to a tuple - a key in the op cache.
-// Op::Make() - a static function, creating the necessary cuDNN descriptor.
-// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor and
+// Op::Make() - a static function, creating all necessary cuDNN descriptors.
+// Op::Clone() - a static function, creating a copy of the op's descriptors with a new cudNN handle.
+// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor(s) and
// the passed arguments.
template <typename Op, typename... Args>
bool Exec(const OpContext& ctx, const typename Op::Param& param, Args&&... args) {
auto key = std::tuple_cat(std::make_tuple(ctx.run_ctx.ctx.dev_id),
Op::MakeKey(param, std::forward<Args>(args)...));
- static std::unordered_map<decltype(key), cudnn_cxx::Descriptor> op_map;
static std::mutex mx;
std::unique_lock<std::mutex> lk(mx);
- auto it = op_map.find(key);
- if (it == op_map.end()) {
- auto op = Op::Make(ctx, param, std::forward<Args>(args)...);
- it = op_map.emplace(key, std::move(op)).first;
- }
+ static std::unordered_multimap<decltype(key), const cudnn_cxx::Descriptor> op_map;
+ auto match_it = [&]() {
+ // Some cuDNN Op implementations require that the thread's cuDNN handle
+ // (used in cudnnBackendExecute()) matches the one used in making the plan.
+ const bool ignore_handles = false;
Review Comment:
The plan is down the road to change `ignore_handles` to be a function of CUDNN_VERSION.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] DickJC123 commented on pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
DickJC123 commented on PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#issuecomment-1199954829
Hey, @mk-61, I think you're the best one to appreciate the intent of this PR and approve it, having supplied the original cuDNN v8 backend support. Could you do a quick review? Thanks!
The existing failure with the windows-gpu job I have diagnosed as being unrelated to this PR. That problem I'm exploring fixes for in separate PR#21107, which I intend to get merged to master, then merged to this PR before final acceptance.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx merged pull request #21104: [BUGFIX] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
ptrendx merged PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a diff in pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
ptrendx commented on code in PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#discussion_r935364141
##########
src/operator/cudnn_ops.h:
##########
@@ -121,25 +128,39 @@ void MaybeLogSelectedPlan(const cudnn_cxx::Descriptor& plan);
// Op::Param - a type, collecting all data, required to create cuDNN descriptor(s), but not needed
// for execution.
// Op::MakeKey() - a static function, which maps its arguments to a tuple - a key in the op cache.
-// Op::Make() - a static function, creating the necessary cuDNN descriptor.
-// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor and
+// Op::Make() - a static function, creating all necessary cuDNN descriptors.
+// Op::Clone() - a static function, creating a copy of the op's descriptors with a new cudNN handle.
+// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor(s) and
// the passed arguments.
template <typename Op, typename... Args>
bool Exec(const OpContext& ctx, const typename Op::Param& param, Args&&... args) {
auto key = std::tuple_cat(std::make_tuple(ctx.run_ctx.ctx.dev_id),
Op::MakeKey(param, std::forward<Args>(args)...));
- static std::unordered_map<decltype(key), cudnn_cxx::Descriptor> op_map;
static std::mutex mx;
std::unique_lock<std::mutex> lk(mx);
- auto it = op_map.find(key);
- if (it == op_map.end()) {
- auto op = Op::Make(ctx, param, std::forward<Args>(args)...);
- it = op_map.emplace(key, std::move(op)).first;
- }
+ static std::unordered_multimap<decltype(key), const cudnn_cxx::Descriptor> op_map;
+ auto match_it = [&]() {
+ // Some cuDNN Op implementations require that the thread's cuDNN handle
+ // (used in cudnnBackendExecute()) matches the one used in making the plan.
+ const bool ignore_handles = false;
Review Comment:
Right now this is always false, right?
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ptrendx commented on a diff in pull request #21104: [BUGFIX] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
ptrendx commented on code in PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#discussion_r937510841
##########
src/operator/cudnn_ops.h:
##########
@@ -121,25 +128,39 @@ void MaybeLogSelectedPlan(const cudnn_cxx::Descriptor& plan);
// Op::Param - a type, collecting all data, required to create cuDNN descriptor(s), but not needed
// for execution.
// Op::MakeKey() - a static function, which maps its arguments to a tuple - a key in the op cache.
-// Op::Make() - a static function, creating the necessary cuDNN descriptor.
-// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor and
+// Op::Make() - a static function, creating all necessary cuDNN descriptors.
+// Op::Clone() - a static function, creating a copy of the op's descriptors with a new cudNN handle.
+// Op::Exec() - a static function, calling cudnnBackendExecute() with the prepared descriptor(s) and
// the passed arguments.
template <typename Op, typename... Args>
bool Exec(const OpContext& ctx, const typename Op::Param& param, Args&&... args) {
auto key = std::tuple_cat(std::make_tuple(ctx.run_ctx.ctx.dev_id),
Op::MakeKey(param, std::forward<Args>(args)...));
- static std::unordered_map<decltype(key), cudnn_cxx::Descriptor> op_map;
static std::mutex mx;
std::unique_lock<std::mutex> lk(mx);
- auto it = op_map.find(key);
- if (it == op_map.end()) {
- auto op = Op::Make(ctx, param, std::forward<Args>(args)...);
- it = op_map.emplace(key, std::move(op)).first;
- }
+ static std::unordered_multimap<decltype(key), const cudnn_cxx::Descriptor> op_map;
+ auto match_it = [&]() {
+ // Some cuDNN Op implementations require that the thread's cuDNN handle
+ // (used in cudnnBackendExecute()) matches the one used in making the plan.
+ const bool ignore_handles = false;
Review Comment:
Ok, makes sense.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#issuecomment-1192144624
Hey @DickJC123 , Thanks for submitting the PR
All tests are already queued to run once. If tests fail, you can trigger one or more tests again with the following commands:
- To trigger all jobs: @mxnet-bot run ci [all]
- To trigger specific jobs: @mxnet-bot run ci [job1, job2]
***
**CI supported jobs**: [clang, unix-cpu, miscellaneous, sanity, centos-cpu, edge, centos-gpu, website, windows-gpu, windows-cpu, unix-gpu]
***
_Note_:
Only following 3 categories can trigger CI :PR Author, MXNet Committer, Jenkins Admin.
All CI tests must pass before the PR can be merged.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] DickJC123 commented on pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
DickJC123 commented on PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#issuecomment-1192222978
Step 1 is complete: with engine 5 reenabled, test_group_conv2d_16c fails in https://jenkins.mxnet-ci.com/blue/organizations/jenkins/mxnet-validation%2Fcentos-gpu/detail/PR-21104/1/pipeline/142.
```
[2022-07-22T05:08:33.227Z] > raise AssertionError(msg)
[2022-07-22T05:08:33.227Z] E AssertionError:
[2022-07-22T05:08:33.227Z] E Items are not equal:
[2022-07-22T05:08:33.227Z] E Error 92150.218750 exceeds tolerance rtol=1.000000e-05, atol=1.000000e-06 (mismatch at least 0.012555%).
[2022-07-22T05:08:33.227Z] E Location of maximum error: (3, 13, 36, 16), a=-0.01365239, b=1.00000000
```
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mk-61 commented on pull request #21104: [BUGFIX] [WIP] Reenable fwd conv engine 5 on test_group_conv2d_16c
Posted by GitBox <gi...@apache.org>.
mk-61 commented on PR #21104:
URL: https://github.com/apache/incubator-mxnet/pull/21104#issuecomment-1199965717
LGTM
--
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.
To unsubscribe, e-mail: commits-unsubscribe@mxnet.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org