[MXNET-11241] Avoid use of troublesome cudnnFind() results when grad_req='add'

[DickJC123]

Description

The problem of issue #11241 arises because cudnnFind() measures convolution algorithm runtimes with an assumed "output blending parameter" beta of 0. However, algorithms may have specialized kernels for the beta==0 case, different and faster than the generalized beta kernels. Should the generalized kernels have issues with the problem size different than the beta==0 kernels, then the algos returned by cudnnFind() might fail when invoked with beta==1 (as it is when the convolution op grad_req='add' argument is present).

The demonstrated problem area involves a large 'c' value of 64K, where for the backprop-to-filter kernel only algo 1 handles the beta==1 case. CudnnFind() was shown to occasionally return algos 0 or 3 as fastest, and both of these return error 8 "execution failed" when run.

The fix is based on the observation that cudnnGet() returns algo 1 for the backprop-to-filter kernel for the troublesome problem sizes. Thus, the fix is to avoid cudnnFind() when grad_req='add' and force use of cudnnGet() instead. The fix maintains the effectiveness of the caching of algo lookups and convolution op instances, so neither cudnnFind() nor cudnnGet() is called repeatedly. Deconvolution was similarly updated with this fix.

Checklist

Essentials

Please feel free to remove inapplicable items for your PR.

• [X ] The PR title starts with [MXNET-$JIRA_ID], where $JIRA_ID refers to the relevant JIRA issue created (except PRs with tiny changes)
• [X ] Changes are complete (i.e. I finished coding on this PR) [1st commit includes test, 2nd the fix]
• [X ] All changes have test coverage:
• Unit tests are added for small changes to verify correctness (e.g. adding a new operator)
• Nightly tests are added for complicated/long-running ones (e.g. changing distributed kvstore)
• Build tests will be added for build configuration changes (e.g. adding a new build option with NCCL)
• [X ] Code is well-documented:
• For user-facing API changes, API doc string has been updated.
• For new C++ functions in header files, their functionalities and arguments are documented.
• For new examples, README.md is added to explain the what the example does, the source of the dataset, expected performance on test set and reference to the original paper if applicable
• Check the API doc at http://mxnet-ci-doc.s3-accelerate.dualstack.amazonaws.com/PR-$PR_ID/$BUILD_ID/index.html
• [X ] To the my best knowledge, examples are either not affected by this change, or have been fixed to be compatible with this change

Changes

• Feature1, tests, (and when applicable, API doc)
• Feature2, tests, (and when applicable, API doc)

Comments

• If this change is a backward incompatible change, why must this change be made.
• Interesting edge cases to note here

[DickJC123]

My second commit did not pass because the CI exposed a flaw in the logic of the fix: while cudnnGet() returned the required algo 1 on Volta, it did not for the Maxwell GPUs that were part of the CI cluster. A third commit narrows the scope of the fix to the observed failure scenarios (grad_req='add' and c>=64K). The fix is to force the problematic backprop-to-filter kernel to use algo 1 in this case. The CI failed on this commit as well, but for seemingly unrelated reasons. Will try an empty commit.

[eric-haibin-lin]

do you mind rebasing on mxnet master and see if CI works?