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

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

ptrendx 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_r332753765
 
 

 ##########
 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:
   You don't need to wait on this event, just recording is fine.

----------------------------------------------------------------
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