diff --git a/src/hip_event.cpp b/src/hip_event.cpp index 7438c7d3d0..73e296dae3 100644 --- a/src/hip_event.cpp +++ b/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(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked + ecd.marker().wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); - return ihipLogStatus(hipSuccess); } } else { diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 1b11b3c67e..1b2f895f7d 100644 --- a/src/hip_hcc.cpp +++ b/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); }; diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index a07ff0e7e7..7ab3fa6de1 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -476,7 +476,7 @@ template class ihipStreamCriticalBase_t : public LockedBase { 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 { 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; }; diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index d70f860f90..7812530dda 100644 --- a/src/hip_stream.cpp +++ b/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(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked - : hc::hcWaitModeActive); + ecd.marker().wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked + : hc::hcWaitModeActive); } else { stream = ihipSyncAndResolveStream(stream); // This will use create_blocking_marker to wait on the specified queue.