Misc fixes to Conv and ConvTranspose CUDA kernels#4281
Merged
hariharans29 merged 4 commits intomasterfrom Jun 29, 2020
Merged
Conversation
hariharans29
changed the title
hariharans29
commented
Jun 19, 2020
| auto b_data = reinterpret_cast<const CudaT*>(B->template Data<T>()); | ||
| CUDNN_RETURN_IF_ERROR(cudnnAddTensor(CudnnHandle(), &alpha, s_.b_tensor, b_data, &alpha, s_.y_tensor, y_data)); | ||
| } | ||
| } |
Member
Author
There was a problem hiding this comment.
Brace from above moved down thus extending holding onto the mutex.
hariharans29
commented
Jun 19, 2020
| ConvTransposeAttributes::Prepare p; | ||
| ORT_RETURN_IF_ERROR(conv_transpose_attrs_.PrepareForCompute(context, has_bias, p, dynamic_padding)); | ||
|
|
||
| // Bail out early if one of the dimensions is zero. |
Member
Author
There was a problem hiding this comment.
Move this block until after s_.y_dims has been cached for future runs
hariharans29
commented
Jun 19, 2020
| s_.algo = perf.algo; | ||
| s_.workspace_bytes = perf.memory; | ||
| } | ||
| } |
Member
Author
There was a problem hiding this comment.
Move this brace down to extend holding onto the mutex longer
hariharans29
commented
Jun 19, 2020
| y_data = reinterpret_cast<CudaT*>(Y->template MutableData<T>()); | ||
|
|
||
| // Bail out early if one of the output dimensions is zero. | ||
| if (Y->Shape().Size() == 0) { |
Member
Author
There was a problem hiding this comment.
if logic to bail out early
skottmckay
reviewed
Jun 26, 2020
hariharans29
changed the title
skottmckay
approved these changes
Jun 29, 2020
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description: There are some subtle fixes to the
ConvandConvTransposeCUDA kernels:In ConvTranspose kernel, extend holding onto the mutex until we are done with all CuDNN calls (Similar to Extend holding onto the mutex in CUDA Conv #3837 which did that for the Conv CUDA kernel)
Fixes corresponding to bailing out too early in the Compute methods if one of the output dim values is 0-
ConvTranpose kernel:
a) There is logic to bail out early if one of the output dimensions is zero, but this logic resides in an
ifblock that only executes for the first run or on a run whose input shape/ filter shape is different from the cached input shape / filter shape from a previous run which meant if the output was empty on one run and the input shape / filter shape remains unchanged on the next run, the logic to bail out early is never hit and will make the CuDNN call resulting in a runtime failure. So fix this.b) If bailing out early in the very first run, cache the
output shapeandfilter dimensionsfor use in subsequent runsConv kernel:
a) If bailing out early in the very first run, cache the
filter dimensionsfor use in subsequent runsMotivation and Context
Fix #4014