remove stream locks where it is safe to do so

src/hip_event.cpp

@@ -191,10 +191,8 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
             ctx->locked_syncDefaultStream(true, true);
             return ihipLogStatus(hipSuccess);
         } else {
-            ecd._stream->locked_eventWaitComplete(
+            ecd.marker().wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
-                ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
                                                                      : hc::hcWaitModeActive);
-
             return ihipLogStatus(hipSuccess);
         }
     } else {

src/hip_hcc.cpp

@@ -316,16 +316,20 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t& crit) {
     tprintf(DB_SYNC, "%s wait for queue-empty..\n", ToString(this).c_str());
 
     crit->_av.wait(waitMode());
-
-    crit->_kernelCnt = 0;
 }
 
 //---
 // Wait for all kernel and data copy commands in this stream to complete.
 void ihipStream_t::locked_wait() {
-    LockedAccessor_StreamCrit_t crit(_criticalData);
+    // create a marker while holding stream lock,
+    // but release lock prior to waiting on the marker
+    hc::completion_future marker;
+    {
+        LockedAccessor_StreamCrit_t crit(_criticalData);
+        marker = crit->_av.create_marker(hc::no_scope);
+    }
 
-    wait(crit);
+    marker.wait(waitMode());
 };
 
 // Causes current stream to wait for specified event to complete:
@@ -340,30 +344,14 @@ void ihipStream_t::locked_streamWaitEvent(ihipEventData_t& ecd) {
 // 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 scrit(_criticalData);
-
     LockedAccessor_EventCrit_t ecrit(event->criticalData());
 
     return (ecrit->_eventData.marker().is_ready());
 }
 
-// Waiting on event can cause HCC to reclaim stream resources - so need to lock the stream.
-void ihipStream_t::locked_eventWaitComplete(hc::completion_future& marker,
-                                            hc::hcWaitMode waitMode) {
-    LockedAccessor_StreamCrit_t crit(_criticalData);
-
-    marker.wait(waitMode);
-}
-
-
 // Create a marker in this stream.
 // Save state in the event so it can track the status of the event.
 hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event) {
-    // Lock the stream to prevent simultaneous access
-    LockedAccessor_StreamCrit_t crit(_criticalData);
-
     auto scopeFlag = hc::accelerator_scope;
     // The env var HIP_EVENT_SYS_RELEASE sets the default,
     // The explicit flags override the env var (if specified)
@@ -375,6 +363,8 @@ hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event) {
         scopeFlag = HIP_EVENT_SYS_RELEASE ? hc::system_scope : hc::accelerator_scope;
     }
 
+    // Lock the stream to prevent simultaneous access
+    LockedAccessor_StreamCrit_t crit(_criticalData);
     return crit->_av.create_marker(scopeFlag);
 };
 

src/hip_hcc_internal.h

@@ -476,7 +476,7 @@ template <typename MUTEX_TYPE>
 class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
    public:
     ihipStreamCriticalBase_t(ihipStream_t* parentStream, hc::accelerator_view av)
-        : _kernelCnt(0), _av(av), _parent(parentStream){};
+        : _av(av), _parent(parentStream){};
 
     ~ihipStreamCriticalBase_t() {}
 
@@ -500,7 +500,6 @@ class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
 
    public:
     ihipStream_t* _parent;
-    uint32_t _kernelCnt;  // Count of inflight kernels in this stream.  Reset at ::wait().
 
     hc::accelerator_view _av;
 
@@ -564,7 +563,6 @@ class ihipStream_t {
     hc::completion_future locked_recordEvent(hipEvent_t event);
 
     bool locked_eventIsReady(hipEvent_t event);
-    void locked_eventWaitComplete(hc::completion_future& marker, hc::hcWaitMode waitMode);
 
     ihipStreamCritical_t& criticalData() { return _criticalData; };
 

src/hip_stream.cpp

@@ -142,9 +142,8 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
                 // conservative wait on host for the specified event to complete:
                 // return _stream->locked_eventWaitComplete(this, waitMode);
                 //
-                ecd._stream->locked_eventWaitComplete(
+                ecd.marker().wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
-                ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
+                                                                         : hc::hcWaitModeActive);
-                                                                     : hc::hcWaitModeActive);
             } else {
                 stream = ihipSyncAndResolveStream(stream);
                 // This will use create_blocking_marker to wait on the specified queue.