From fe85edad3ff50c190fb1fa606851e24b4d1f39f4 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 1 Sep 2016 13:59:55 -0500
Subject: [PATCH] Use create_blocking_marker for WaitEvent implementation

Change-Id: Ib3113f69a14e48b9fe0558d7b455148e478d8eed


[ROCm/clr commit: 9a99000a62f70556c9f43f488c0f0fee360e19da]
---
 projects/clr/hipamd/src/hip_event.cpp  | 19 +++++++--------
 projects/clr/hipamd/src/hip_hcc.cpp    | 10 ++++++++
 projects/clr/hipamd/src/hip_stream.cpp | 32 +++++++++++++++++++-------
 3 files changed, 42 insertions(+), 19 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp
index 33ec0bc897..ecc6038b82 100644
--- a/projects/clr/hipamd/src/hip_event.cpp
+++ b/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);
diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index ca650e5283..1854e3d8c7 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/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.
diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp
index 5d900fcaeb..f400e27799 100644
--- a/projects/clr/hipamd/src/hip_stream.cpp
+++ b/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);
 };