diff --git a/hipamd/src/hip_device_runtime.cpp b/hipamd/src/hip_device_runtime.cpp index 19bed5ef83..5376ae0faf 100644 --- a/hipamd/src/hip_device_runtime.cpp +++ b/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(stream)) == true) { HIP_RETURN(hipErrorStreamCaptureUnsupported); } diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index deeea225ad..f9a699e130 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -29,7 +29,7 @@ std::vector g_captureStreams; amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"}; -static amd::Monitor g_streamSetLock{"StreamCaptureset"}; +amd::Monitor g_streamSetLock{"StreamCaptureset"}; std::unordered_set g_allCapturingStreams; inline hipError_t ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph, diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index ca924bd5f8..2983a934c7 100644 --- a/hipamd/src/hip_internal.hpp +++ b/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 g_captureStreams; extern amd::Monitor g_captureStreamsLock; +extern amd::Monitor g_streamSetLock; extern std::unordered_set g_allCapturingStreams; #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 45ddd1665b..8a240a857f 100644 --- a/hipamd/src/hip_stream.cpp +++ b/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(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));