From bd434502db5eeec0b7c19fe6877bb230ae114bcf Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 26 Aug 2017 14:39:14 +0000 Subject: [PATCH] Lock streams when waiting on event completion or querying event safety. [ROCm/hip commit: 6ff74d0e977a403a7469e5c023c67a53b9c7878d] --- projects/hip/src/hip_event.cpp | 23 ++++++++++++++++------- projects/hip/src/hip_hcc.cpp | 26 +++++++++++++++++++++++--- projects/hip/src/hip_hcc_internal.h | 15 ++++++++++++--- projects/hip/src/hip_stream.cpp | 2 +- 4 files changed, 52 insertions(+), 14 deletions(-) diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index 3a8f1ab611..d1ee37a45e 100644 --- a/projects/hip/src/hip_event.cpp +++ b/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 (_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); - - assert (event->_marker.is_ready()); + event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); 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); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index f7082e7e9e..dc72714e3e 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/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)); }; //============================================================================= diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index af5e7a121b..56ca37b3e2 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/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); diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index f0e9283201..51aeb01412 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/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.