SWDEV-387810 - Fixes to verify if the current stream matches the capture stream

- invalid to synchronize or query the execution status of a capturing stream but non-capturing streams can.

- If any local thread has an ongoing or concurrent capture sequence initiated with hipStreamCaptureModeGlobal, it is prohibited from unsafe calls

Change-Id: Ifa641e807216d3b7c3e8c2fb1be2f7a50bd641df

hipamd/src/hip_device_runtime.cpp

@@ -518,7 +518,7 @@ hipError_t hipDeviceSynchronize ( void ) {
     HIP_RETURN(hipErrorOutOfMemory);
   }
 
-  if (hip::Stream::StreamCaptureOngoing() == true) {
+  if (hip::Stream::StreamCaptureOngoing(reinterpret_cast<hipStream_t>(stream)) == true) {
     HIP_RETURN(hipErrorStreamCaptureUnsupported);
   }
 

hipamd/src/hip_graph.cpp

@@ -29,7 +29,7 @@
 
 std::vector<hip::Stream*> g_captureStreams;
 amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"};
-static amd::Monitor g_streamSetLock{"StreamCaptureset"};
+amd::Monitor g_streamSetLock{"StreamCaptureset"};
 std::unordered_set<hip::Stream*> g_allCapturingStreams;
 
 inline hipError_t ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph,

hipamd/src/hip_internal.hpp

@@ -300,7 +300,7 @@ namespace hip {
     static void destroyAllStreams(int deviceId);
 
     /// Check Stream Capture status to make sure it is done
-    static bool StreamCaptureOngoing(void);
+    static bool StreamCaptureOngoing(hipStream_t hStream);
 
     /// Returns capture status of the current stream
     hipStreamCaptureStatus GetCaptureStatus() const { return captureStatus_; }
@@ -574,5 +574,6 @@ constexpr bool kMarkerDisableFlush = true;   //!< Avoids command batch flush in
 
 extern std::vector<hip::Stream*> g_captureStreams;
 extern amd::Monitor g_captureStreamsLock;
+extern amd::Monitor g_streamSetLock;
 extern std::unordered_set<hip::Stream*> g_allCapturingStreams;
 #endif // HIP_SRC_HIP_INTERNAL_H

hipamd/src/hip_stream.cpp

@@ -156,8 +156,18 @@ void Stream::destroyAllStreams(int deviceId) {
   }
 }
 
-bool Stream::StreamCaptureOngoing(void) {
-  return (g_allCapturingStreams.empty() == true) ? false : true;
+bool Stream::StreamCaptureOngoing(hipStream_t hStream) {
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(hStream);
+  // If any local thread has an ongoing or concurrent capture sequence initiated
+  // with hipStreamCaptureModeGlobal, it is prohibited from unsafe calls
+  if (s != nullptr && s->GetCaptureMode() == hipStreamCaptureModeGlobal) {
+    amd::ScopedLock lock(g_captureStreamsLock);
+    return (g_captureStreams.empty() == true) ? false : true;
+  }
+  else {
+    amd::ScopedLock lock(g_streamSetLock);
+    return (g_allCapturingStreams.find(s) == g_allCapturingStreams.end() ? false : true);
+  }
 }
 
 bool Stream::existsActiveStreamForDevice(hip::Device* device) {
@@ -416,7 +426,7 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) {
   }
   if (stream != nullptr) {
     // If still capturing return error
-    if (hip::Stream::StreamCaptureOngoing() == true) {
+    if (hip::Stream::StreamCaptureOngoing(stream) == true) {
       HIP_RETURN(hipErrorStreamCaptureUnsupported);
     }
   }
@@ -526,7 +536,7 @@ hipError_t hipStreamQuery_common(hipStream_t stream) {
   }
   if (stream != nullptr) {
     // If still capturing return error
-    if (hip::Stream::StreamCaptureOngoing() == true) {
+    if (hip::Stream::StreamCaptureOngoing(stream) == true) {
       HIP_RETURN(hipErrorStreamCaptureUnsupported);
     }
   }
@@ -657,7 +667,7 @@ hipError_t hipLaunchHostFunc_spt(hipStream_t stream, hipHostFn_t fn, void* userD
 // ================================================================================================
 hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData) {
   HIP_INIT_API(hipLaunchHostFunc, stream, fn, userData);
-  if (stream == nullptr && (hip::Stream::StreamCaptureOngoing() == true)) {
+  if (stream == nullptr && (hip::Stream::StreamCaptureOngoing(stream) == true)) {
     HIP_RETURN(hipErrorStreamCaptureImplicit);
   }
   HIP_RETURN(hipLaunchHostFunc_common(stream, fn, userData));