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