diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index fbaf5cc463..c11a47b341 100644 --- a/hipamd/src/hip_event.cpp +++ b/hipamd/src/hip_event.cpp @@ -53,14 +53,19 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, ihip void ihipEvent_t::setTimestamp() { + bool isReady0 = _marker.is_ready(); + bool isReady1; + int val = 0; if (_state == hipEventStatusRecorded) { // already recorded, done: return; } else { // TODO - use completion-future functions to obtain ticks and timestamps: hsa_signal_t *sig = static_cast (_marker.get_native_handle()); + isReady1 = _marker.is_ready(); if (sig) { - if (hsa_signal_load_acquire(*sig) == 0) { + val = hsa_signal_load_acquire(*sig); + if (val == 0) { if ((_type == hipEventTypeIndependent) || (_type == hipEventTypeStopCommand)) { _timestamp = _marker.get_end_tick(); @@ -75,6 +80,10 @@ void ihipEvent_t::setTimestamp() } } } + + if (_state != hipEventStatusRecorded) { + printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1); + } } @@ -118,11 +127,11 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) event->_stream = stream; - if (HIP_SYNC_NULL_STREAM && stream == NULL) { + if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) { // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0 - // If stream == NULL, wait on all queues. + // If default stream , then wait on all queues. ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); ctx->locked_syncDefaultStream(true, true); @@ -167,7 +176,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) } else if (event->_state == hipEventStatusCreated ) { // Created but not actually recorded on any device: return ihipLogStatus(hipSuccess); - } else if (HIP_SYNC_NULL_STREAM && (event->_stream == NULL)) { + } else if (HIP_SYNC_NULL_STREAM && (event->_stream->isDefaultStream() )) { auto *ctx = ihipGetTlsDefaultCtx(); // TODO-HIP_SYNC_NULL_STREAM - can remove this code ctx->locked_syncDefaultStream(true, true); @@ -175,6 +184,8 @@ hipError_t hipEventSynchronize(hipEvent_t event) } else { event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); + assert (event->_marker.is_ready()); + return ihipLogStatus(hipSuccess); } } else { diff --git a/hipamd/src/hip_hcc_internal.h b/hipamd/src/hip_hcc_internal.h index 278f52dc51..94ad4f9340 100644 --- a/hipamd/src/hip_hcc_internal.h +++ b/hipamd/src/hip_hcc_internal.h @@ -538,10 +538,12 @@ public: const ihipDevice_t * getDevice() const; ihipCtx_t * getCtx() const; + bool isDefaultStream() const { return _id == 0; }; + public: //--- //Public member vars - these are set at initialization and never change: - SeqNum_t _id; // monotonic sequence ID + SeqNum_t _id; // monotonic sequence ID. 0 is the default stream. unsigned _flags; @@ -560,6 +562,7 @@ private: void addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes); + public: // TODO - move private // Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t ihipStreamCritical_t _criticalData; diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index b4a0740b96..9f1228d6f7 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -146,7 +146,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream) hipError_t e = hipSuccess; - if (stream == NULL) { + if (stream == hipStreamNull) { ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/); } else { @@ -198,7 +198,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) if (flags == NULL) { return ihipLogStatus(hipErrorInvalidValue); - } else if (stream == NULL) { + } else if (stream == hipStreamNull) { return ihipLogStatus(hipErrorInvalidResourceHandle); } else { *flags = stream->_flags;