From dc97f3719e5abf748097819ba222c5f99df5dd4c Mon Sep 17 00:00:00 2001
From: sdashmiz <shadi.dashmiz@amd.com>
Date: Fri, 11 Nov 2022 11:41:00 -0500
Subject: [PATCH] SWDEV-366974 - add proper error

- Null stream correct error when blocing stream is being captured

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I0d902a4f4ef02736a6db60ba7bffa7c20011bc84


[ROCm/clr commit: c5e7b56463b4920f5fce34baa143662170b2a176]
---
 projects/clr/hipamd/src/hip_graph.cpp    |  3 +++
 projects/clr/hipamd/src/hip_internal.hpp |  3 +++
 projects/clr/hipamd/src/hip_stream.cpp   | 10 ++++++++++
 3 files changed, 16 insertions(+)

diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp
index 001de82924..caa9183c0c 100644
--- a/projects/clr/hipamd/src/hip_graph.cpp
+++ b/projects/clr/hipamd/src/hip_graph.cpp
@@ -807,6 +807,9 @@ hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatu
   if (pCaptureStatus == nullptr || !hip::isValid(stream)) {
     return hipErrorInvalidValue;
   }
+  if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) {
+    return hipErrorStreamCaptureImplicit;
+  }
   if (stream == nullptr) {
     *pCaptureStatus = hipStreamCaptureStatusNone;
   } else {
diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp
index b7795b9842..60d584a9c9 100644
--- a/projects/clr/hipamd/src/hip_internal.hpp
+++ b/projects/clr/hipamd/src/hip_internal.hpp
@@ -295,6 +295,9 @@ namespace hip {
     /// Sync all non-blocking streams
     static void syncNonBlockingStreams(int deviceId);
 
+    /// Check whether any blocking stream running
+    static bool StreamCaptureBlocking();
+
     /// Destroy all streams on a given device
     static void destroyAllStreams(int deviceId);
 
diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp
index 19fb6fc0ad..d65f738dda 100644
--- a/projects/clr/hipamd/src/hip_stream.cpp
+++ b/projects/clr/hipamd/src/hip_stream.cpp
@@ -180,6 +180,16 @@ void Stream::syncNonBlockingStreams(int deviceId) {
   }
 }
 
+bool Stream::StreamCaptureBlocking() {
+  amd::ScopedLock lock(streamSetLock);
+  for (auto& it : streamSet) {
+    if (it->GetCaptureStatus() == hipStreamCaptureStatusActive && it->Flags() != hipStreamNonBlocking) {
+      return true;
+    }
+  }
+  return false;
+}
+
 void Stream::destroyAllStreams(int deviceId) {
   std::vector<Stream*> toBeDeleted;
   {