DickJC123 commented on a change in pull request #16391: cuDNN non-persistant
bidirectional RNN dgrad sync fix
URL: https://github.com/apache/incubator-mxnet/pull/16391#discussion_r332763755
##########
File path: src/operator/rnn-inl.h
##########
@@ -1484,6 +1493,20 @@ class RNNOp {
}
#endif // MXNET_USE_CUDNN == 1 && defined(__CUDACC__)
}
+
+#if MXNET_USE_CUDNN == 1 && defined(__CUDACC__)
+ // cuDNN versions up to and including v7.6.4 did not sync a last dgrad
kernel back to the main
+ // cudnn handle's stream (non-persistant algo, bidirectional only). This
could result in silent
+ // non-determinstic failures with very low probability, seen more often when
wgrad is bypassed.
+ inline void SyncDgrad() {
+ if (CUDNN_VERSION <= 7604 && dgrad_sync_needed_) {
+ // Without blocking the CPU, create a synchronization point of all
current GPU activity.
+ CUDA_CALL(cudaEventRecord(dgrad_sync_event_, cudaStreamLegacy));
+ CUDA_CALL(cudaStreamWaitEvent(cudaStreamLegacy, dgrad_sync_event_, 0));
Review comment:
I've simplified this as you directed, but put in a comment to say that the
wait is not needed in this case.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
[email protected]
With regards,
Apache Git Services
