SWDEV-326798 - Changes in stream sync behavior

Change-Id: If6d0b3876a9bf197c7e49273eaa5ca5bfae46d0b


[ROCm/clr commit: 51adec2730]

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<hip::Stream*>(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;
 }
 
 };

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_;
 }
 

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

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) {