SWDEV-326798 - [ABI Break] Fixes null stream sync behavior

Change-Id: I84ad21f61779145c198dc68bf0fe02a57d34bf64


[ROCm/clr commit: 3f88fe850b]

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

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

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