From 1ef9123b54ece6dd572e49bd3380cbd91e7c4b92 Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Wed, 14 Aug 2024 14:59:36 +0000
Subject: [PATCH] SWDEV-478776 - Fix segfault with streamsync using
 hipStreamLegacy

Change-Id: Ifb412d0bcfa33bc1130b47b757ee276ca9bc1c3a


[ROCm/clr commit: 8ee065c5bd1d536bc6dea60a378e03d460e18f25]
---
 projects/clr/hipamd/src/hip_stream.cpp | 14 +++++++++-----
 1 file changed, 9 insertions(+), 5 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp
index 233db0b05f..bfd1c965df 100644
--- a/projects/clr/hipamd/src/hip_stream.cpp
+++ b/projects/clr/hipamd/src/hip_stream.cpp
@@ -136,8 +136,12 @@ bool Stream::StreamCaptureBlocking() {
 }
 
 bool Stream::StreamCaptureOngoing(hipStream_t hStream) {
+  if (hStream == nullptr || hStream == hipStreamLegacy) {
+    return false;
+  }
+
   hip::Stream* s = reinterpret_cast<hip::Stream*>(hStream);
-  if (s != nullptr && s->GetCaptureStatus() == hipStreamCaptureStatusNone) {
+  if (s->GetCaptureStatus() == hipStreamCaptureStatusNone) {
     // If current thread is capturing in relaxed mode
     if (hip::tls.stream_capture_mode_ == hipStreamCaptureModeRelaxed) {
       return false;
@@ -158,10 +162,10 @@ bool Stream::StreamCaptureOngoing(hipStream_t hStream) {
       return true;
     }
     return false;
-  } else if (s != nullptr && s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
+  } else if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
     s->SetCaptureStatus(hipStreamCaptureStatusInvalidated);
     return true;
-  } else if (s != nullptr && s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) {
+  } else if (s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) {
     return true;
   }
   return false;
@@ -346,13 +350,13 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) {
   if (!hip::isValid(stream)) {
     HIP_RETURN(hipErrorContextIsDestroyed);
   }
-  if (stream != nullptr) {
+  if (stream != nullptr && stream != hipStreamLegacy) {
     // If still capturing return error
     if (hip::Stream::StreamCaptureOngoing(stream) == true) {
       HIP_RETURN(hipErrorStreamCaptureUnsupported);
     }
   }
-  bool wait = (stream == nullptr) ? true : false;
+  bool wait = (stream == nullptr || stream == hipStreamLegacy) ? true : false;
   constexpr bool kDontWaitForCpu = false;
 
   auto hip_stream = hip::getStream(stream, wait);