Lock streams when waiting on event completion or querying event safety.

[ROCm/hip commit: 6ff74d0e97]

projects/hip/src/hip_event.cpp

@@ -55,13 +55,13 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf,
 
 void ihipEvent_t::refereshEventStatus()
 {
-    bool isReady0 = _marker.is_ready();
+    bool isReady0 = locked_isReady();
     bool isReady1;
     int val = 0;
     if (_state == hipEventStatusRecording) {
         // TODO - use completion-future functions to obtain ticks and timestamps:
         hsa_signal_t *sig  = static_cast<hsa_signal_t*> (_marker.get_native_handle());
-        isReady1 = _marker.is_ready();
+        isReady1 = locked_isReady();
         if (sig) {
             val = hsa_signal_load_acquire(*sig);
             if (val == 0) {
@@ -86,6 +86,17 @@ void ihipEvent_t::refereshEventStatus()
 }
 
 
+bool ihipEvent_t::locked_isReady()
+{
+    return _stream->locked_eventIsReady(this);
+}
+
+void ihipEvent_t::locked_waitComplete(hc::hcWaitMode waitMode)
+{
+    return _stream->locked_eventWaitComplete(this, waitMode);
+}
+
+
 hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
 {
     hipError_t e = hipSuccess;
@@ -127,7 +138,7 @@ hipError_t hipEventCreate(hipEvent_t* event)
 
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
 {
-    HIP_INIT_SPECIAL_API(TRACE_QUERY, event, stream);
+    HIP_INIT_SPECIAL_API(TRACE_SYNC, event, stream);
 
     if (event && event->_state != hipEventStatusUnitialized)   {
         stream = ihipSyncAndResolveStream(stream);
@@ -192,9 +203,7 @@ hipError_t hipEventSynchronize(hipEvent_t event)
             ctx->locked_syncDefaultStream(true, true);
             return ihipLogStatus(hipSuccess);
         } else {
-            event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
+            event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
-
-            assert (event->_marker.is_ready());
 
             return ihipLogStatus(hipSuccess);
         }
@@ -259,7 +268,7 @@ hipError_t hipEventQuery(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
 
-    if ((event->_state == hipEventStatusRecording) && (!event->_marker.is_ready())) {
+    if ((event->_state == hipEventStatusRecording) && !event->locked_isReady()) {
         return ihipLogStatus(hipErrorNotReady);
     } else {
         return ihipLogStatus(hipSuccess);

projects/hip/src/hip_hcc.cpp

@@ -328,14 +328,34 @@ void ihipStream_t::locked_wait()
 
 // Causes current stream to wait for specified event to complete:
 // Note this does not provide any kind of host serialization.
-void ihipStream_t::locked_waitEvent(hipEvent_t event)
+void ihipStream_t::locked_streamWaitEvent(hipEvent_t event)
 {
     LockedAccessor_StreamCrit_t crit(_criticalData);
 
 
-    crit->_av.create_blocking_marker(event->_marker, hc::accelerator_scope);
+    crit->_av.create_blocking_marker(event->marker(), hc::accelerator_scope);
 }
 
+
+// Causes current stream to wait for specified event to complete:
+// Note this does not provide any kind of host serialization.
+bool ihipStream_t::locked_eventIsReady(hipEvent_t event)
+{
+    // Event query that returns "Complete" may cause HCC to manipulate
+    // internal queue state so lock the stream's queue here.
+    LockedAccessor_StreamCrit_t crit(_criticalData);
+
+    return (event->marker().is_ready());
+}
+
+void ihipStream_t::locked_eventWaitComplete(hipEvent_t event, hc::hcWaitMode waitMode)
+{
+    LockedAccessor_StreamCrit_t crit(_criticalData);
+
+    event->marker().wait(waitMode);
+}
+
+
 // Create a marker in this stream.
 // Save state in the event so it can track the status of the event.
 void ihipStream_t::locked_recordEvent(hipEvent_t event)
@@ -354,7 +374,7 @@ void ihipStream_t::locked_recordEvent(hipEvent_t event)
         scopeFlag = HIP_EVENT_SYS_RELEASE ? hc::system_scope : hc::accelerator_scope;
     }
 
-    event->_marker = crit->_av.create_marker(scopeFlag);
+    event->marker(crit->_av.create_marker(scopeFlag));
 };
 
 //=============================================================================

projects/hip/src/hip_hcc_internal.h

@@ -517,9 +517,12 @@ public:
 
     hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); };
 
-    void                 locked_waitEvent(hipEvent_t event);
+    void                 locked_streamWaitEvent(hipEvent_t event);
     void                 locked_recordEvent(hipEvent_t event);
 
+    bool                 locked_eventIsReady(hipEvent_t event);
+    void                 locked_eventWaitComplete(hipEvent_t event, hc::hcWaitMode waitMode);
+
     ihipStreamCritical_t  &criticalData() { return _criticalData; };
 
     //---
@@ -608,18 +611,24 @@ public:
     ihipEvent_t(unsigned flags);
     void attachToCompletionFuture(const hc::completion_future *cf, hipStream_t stream, ihipEventType_t eventType);
     void refereshEventStatus();
+    hc::completion_future & marker() { return _marker; }
+    void marker(hc::completion_future cf) { _marker = cf; };
+
+    bool locked_isReady();
+    void locked_waitComplete(hc::hcWaitMode waitMode);
+
     uint64_t timestamp() const { return _timestamp; } ;
     ihipEventType_t type() const { return _type; };
 
 public:
     hipEventStatus_t      _state;
 
-    hipStream_t           _stream;  // Stream where the event is recorded, or NULL if all streams.
+    hipStream_t           _stream;  // Stream where the event is recorded.  Null stream is resolved to actual stream when recorded
     unsigned              _flags;
 
-    hc::completion_future _marker;
 
 private:
+    hc::completion_future _marker;
     ihipEventType_t       _type;
     uint64_t              _timestamp;  // store timestamp, may be set on host or by marker.
 friend hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream);

projects/hip/src/hip_stream.cpp

@@ -96,7 +96,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
         if (stream != hipStreamNull) {
 
             // This will user create_blocking_marker to wait on the specified queue.
-            stream->locked_waitEvent(event);
+            stream->locked_streamWaitEvent(event);
 
         } else {
             // TODO-hcc Convert to use create_blocking_marker(...) functionality.