Use create_blocking_marker for WaitEvent implementation

Change-Id: Ib3113f69a14e48b9fe0558d7b455148e478d8eed


[ROCm/clr commit: 9a99000a62]

projects/clr/hipamd/src/hip_event.cpp

@@ -40,7 +40,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
         eh->_flags  = flags;
         eh->_timestamp  = 0;
         eh->_copySeqId  = 0;
-        *event = eh;
+        *event = eh; // TODO - allocat the event directly, no copy needed.
     } else {
         e = hipErrorInvalidValue;
     }
@@ -121,21 +121,19 @@ hipError_t hipEventSynchronize(hipEvent_t event)
 {
     HIP_INIT_API(event);
 
-    ihipEvent_t *eh = event;
-
-    if (eh) {
-        if (eh->_state == hipEventStatusUnitialized) {
+    if (event) {
+        if (event->_state == hipEventStatusUnitialized) {
             return ihipLogStatus(hipErrorInvalidResourceHandle);
-        } else if (eh->_state == hipEventStatusCreated ) {
+        } else if (event->_state == hipEventStatusCreated ) {
             // Created but not actually recorded on any device:
             return ihipLogStatus(hipSuccess);
-        } else if (eh->_stream == NULL) {
+        } else if (event->_stream == NULL) {
             auto *ctx = ihipGetTlsDefaultCtx();
             ctx->locked_syncDefaultStream(true);
             return ihipLogStatus(hipSuccess);
         } else {
-            eh->_marker.wait((eh->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
-            eh->_stream->locked_reclaimSignals(eh->_copySeqId);
+            event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
+            event->_stream->locked_reclaimSignals(event->_copySeqId);
 
             return ihipLogStatus(hipSuccess);
         }
@@ -195,12 +193,11 @@ hipError_t hipEventQuery(hipEvent_t event)
 {
     HIP_INIT_API(event);
 
-    ihipEvent_t *eh = event;
 
     // TODO-stream - need to read state of signal here:  The event may have become ready after recording..
     // TODO-HCC - use get_hsa_signal here.
 
-    if (eh->_state == hipEventStatusRecording) {
+    if (event->_state == hipEventStatusRecording) {
         return ihipLogStatus(hipErrorNotReady);
     } else {
         return ihipLogStatus(hipSuccess);

projects/clr/hipamd/src/hip_hcc.cpp

@@ -273,6 +273,16 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty)
 
 };
 
+#if USE_AV_COPY
+// Causes current stream to wait for specified event to complete:
+void ihipStream_t::locked_waitEvent(hipEvent_t event)
+{
+    LockedAccessor_StreamCrit_t crit(_criticalData);
+
+    // TODO - check state of event here:
+    crit->_av.create_blocking_marker(event->_marker);
+}
+#endif
 
 // Create a marker in this stream.
 // Save state in the event so it can track the status of the event.

projects/clr/hipamd/src/hip_stream.cpp

@@ -74,27 +74,43 @@ hipError_t hipStreamCreate(hipStream_t *stream)
 }
 
 
+#ifndef USE_AV_COPY
 //---
 /**
  * @bug This function conservatively waits for all work in the specified stream to complete.
  */
+#endif
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
 {
     HIP_INIT_API(stream, event, flags);
 
     hipError_t e = hipSuccess;
 
+    if (event == nullptr) {
+        e = hipErrorInvalidResourceHandle;
+
+    } else if (event->_state != hipEventStatusUnitialized) {
+
+        bool fastWait = false;
+
 #ifdef USE_AV_COPY
-     printf ("USE_AV_COPY\n");
+        if (stream != hipStreamNull) {
+            printf ("HIP: wait locked stream\n");
+            stream->locked_waitEvent(event);
+
+            fastWait = true; // don't use the slow host-side synchronization.
+        }
+        // TODO - clean up if/else logic when USE_AV_COPY enabled.
 #endif
 
-    {
-        // TODO-hcc Convert to use create_blocking_marker(...) functionality.
-        // Currently we have a super-conservative version of this - block on host, and drain the queue.
-        // This should create a barrier packet in the target queue.
-        stream->locked_wait();
-        e = hipSuccess;
-    }
+        if (!fastWait) {
+            // TODO-hcc Convert to use create_blocking_marker(...) functionality.
+            // Currently we have a super-conservative version of this - block on host, and drain the queue.
+            // This should create a barrier packet in the target queue.
+            stream->locked_wait();
+            e = hipSuccess;
+        }
+    } // else event not recorded, return immediately and don't create marker.
 
     return ihipLogStatus(e);
 };