[GitHub] [incubator-mxnet] DickJC123 commented on a change in pull request #16391: cuDNN non-persistant bidirectional RNN dgrad sync fix

[GitBox <gi...@apache.org>]

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:
users@infra.apache.org


With regards,
Apache Git Services