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:
[email protected]
With regards,
Apache Git Services
