diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index f639d4ff66..dfcd189c67 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -91,12 +91,12 @@ void setCurrentDevice(unsigned int index) { amd::Os::setPreferredNumaNode(preferredNumaNode); } -hip::Stream* getStream(hipStream_t stream) { +hip::Stream* getStream(hipStream_t stream, bool wait) { if (stream == nullptr) { - return getNullStream(); + return getNullStream(wait); } else { hip::Stream* hip_stream = reinterpret_cast(stream); - if (!(hip_stream->Flags() & hipStreamNonBlocking)) { + if (wait && !(hip_stream->Flags() & hipStreamNonBlocking)) { constexpr bool WaitNullStreamOnly = true; iHipWaitActiveStreams(hip_stream, WaitNullStreamOnly); } @@ -131,9 +131,10 @@ int getDeviceID(amd::Context& ctx) { } // ================================================================================================ -hip::Stream* getNullStream() { +hip::Stream* getNullStream(bool wait) { Device* device = getCurrentDevice(); - return device ? device->NullStream() : nullptr; + constexpr bool kSkipAlloc = false; + return device ? device->NullStream(kSkipAlloc, wait) : nullptr; } }; diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index 8782b6c35e..82fcee6201 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -26,7 +26,7 @@ namespace hip { // ================================================================================================ -hip::Stream* Device::NullStream(bool skip_alloc) { +hip::Stream* Device::NullStream(bool skip_alloc, bool wait) { if (null_stream_ == nullptr && !skip_alloc) { null_stream_ = new Stream(this, Stream::Priority::Normal, 0, true); } @@ -34,8 +34,10 @@ hip::Stream* Device::NullStream(bool skip_alloc) { if (null_stream_ == nullptr) { return nullptr; } - // Wait for all active streams before executing commands on the default - iHipWaitActiveStreams(null_stream_); + if (wait == true) { + // Wait for all active streams before executing commands on the default + iHipWaitActiveStreams(null_stream_); + } return null_stream_; } diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 49a8832650..e92eaeb768 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -451,10 +451,9 @@ namespace hip { void setFlags(unsigned int flags) { flags_ = flags; } void Reset(); - hip::Stream* NullStream(bool skip_alloc = false); + hip::Stream* NullStream(bool skip_alloc = false, bool wait = true); Stream* GetNullStream(); - bool GetActiveStatus() { amd::ScopedLock lock(lock_); if (isActive_) return true; @@ -528,11 +527,11 @@ namespace hip { /// Get ROCclr queue associated with hipStream /// Note: This follows the CUDA spec to sync with default streams /// and Blocking streams - extern hip::Stream* getStream(hipStream_t stream); + extern hip::Stream* getStream(hipStream_t stream, bool wait = true); /// Get default stream associated with the ROCclr context extern hip::Stream* getNullStream(amd::Context&); /// Get default stream of the thread - extern hip::Stream* getNullStream(); + extern hip::Stream* getNullStream(bool wait = true); /// Get device ID associated with the ROCclr context int getDeviceID(amd::Context& ctx); /// Check if stream is valid diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 1342ac7270..28550c48c0 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -420,8 +420,9 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) { HIP_RETURN(hipErrorStreamCaptureUnsupported); } } + bool wait = (stream == nullptr) ? true : false; // Wait for the current host queue - hip::getStream(stream)->finish(); + hip::getStream(stream, wait)->finish(); return hipSuccess; } @@ -530,7 +531,9 @@ hipError_t hipStreamQuery_common(hipStream_t stream) { HIP_RETURN(hipErrorStreamCaptureUnsupported); } } - hip::Stream* hip_stream = hip::getStream(stream); + + bool wait = (stream == nullptr) ? true : false; + hip::Stream* hip_stream = hip::getStream(stream, wait); amd::Command* command = hip_stream->getLastQueuedCommand(true); if (command == nullptr) {