From 5d160eb7ee27dfa691cb8764206a8d0535ee160a Mon Sep 17 00:00:00 2001
From: Sourabh Betigeri <sourabh.betigeri@amd.com>
Date: Wed, 26 Apr 2023 20:27:11 +0000
Subject: [PATCH] SWDEV-326798 - [ABI Break] Fixes null stream sync behavior

Change-Id: I84ad21f61779145c198dc68bf0fe02a57d34bf64


[ROCm/clr commit: 3f88fe850be04b6784c204e6beacf772174a99a2]
---
 projects/clr/hipamd/src/hip_context.cpp  | 4 ++--
 projects/clr/hipamd/src/hip_internal.hpp | 2 +-
 projects/clr/hipamd/src/hip_stream.cpp   | 6 ++++--
 3 files changed, 7 insertions(+), 5 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp
index a9ef183631..2be3b77f5f 100644
--- a/projects/clr/hipamd/src/hip_context.cpp
+++ b/projects/clr/hipamd/src/hip_context.cpp
@@ -100,12 +100,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();
   } 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);
     }
diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp
index 1cecf0547f..2cef6902ed 100644
--- a/projects/clr/hipamd/src/hip_internal.hpp
+++ b/projects/clr/hipamd/src/hip_internal.hpp
@@ -552,7 +552,7 @@ 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
diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp
index 8ee1ce9e9e..138fa018fc 100644
--- a/projects/clr/hipamd/src/hip_stream.cpp
+++ b/projects/clr/hipamd/src/hip_stream.cpp
@@ -441,8 +441,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;
 }
 
@@ -572,7 +573,8 @@ 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) {