Merge pull request #1083 from dkolsen-pgi/bug/cudaStreamWaitEvent

Add third argument to calls to cudaStreamWaitEvent

Author: ericniebler

include/nvexec/stream/bulk.cuh

@@ -202,7 +202,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
 
                 if (begin < end) {
                   cudaSetDevice(dev);
-                  cudaStreamWaitEvent(stream, op_state.ready_to_launch_);
+                  cudaStreamWaitEvent(stream, op_state.ready_to_launch_, 0);
                   kernel<block_threads, As&...>
                     <<<grid_blocks, block_threads, 0, stream>>>(begin, end, self.f_, as...);
                   cudaEventRecord(op_state.ready_to_complete_[dev], op_state.streams_[dev]);
@@ -225,7 +225,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
 
             for (int dev = 0; dev < op_state.num_devices_; dev++) {
               if (dev != op_state.current_device_) {
-                cudaStreamWaitEvent(baseline_stream, op_state.ready_to_complete_[dev]);
+                cudaStreamWaitEvent(baseline_stream, op_state.ready_to_complete_[dev], 0);
               }
             }
           }

include/nvexec/stream/ensure_started.cuh

@@ -246,7 +246,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
           if (status == cudaSuccess) {
             if constexpr (stream_sender<Sender, env_t>) {
               status = STDEXEC_DBG_ERR(
-                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_));
+                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_, 0));
             }
 
             visit(

include/nvexec/stream/split.cuh

@@ -236,7 +236,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
           if (status == cudaSuccess) {
             if constexpr (stream_sender<Sender, env_t>) {
               status = STDEXEC_DBG_ERR(
-                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_));
+                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_, 0));
             }
 
             visit(

include/nvexec/stream/when_all.cuh

@@ -259,7 +259,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
 
               for (int i = 0; i < sizeof...(SenderIds); i++) {
                 if (status_ == cudaSuccess) {
-                  status_ = STDEXEC_DBG_ERR(cudaStreamWaitEvent(stream, events_[i]));
+                  status_ = STDEXEC_DBG_ERR(cudaStreamWaitEvent(stream, events_[i], 0));
                 }
               }
             } else {