From ae296c8fadf0748f757c41695b8dd6e73ea00c18 Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Thu, 25 Jan 2024 10:32:08 +0000 Subject: [PATCH] SWDEV-443567 - SWDEV-436126 - Fix Prohibited and Unhandled Operations during capture => hipDeviceSynchronize is not allowed during capture. => hipEventSynchronize during capture should return hipErrorCapturedEvent error => hipEventQuery during capture should return hipErrorCapturedEvent error hipStreamSynchronize, hipEventSynchronize, hipStreamWaitEvent, hipStreamQuery For Side Stream(Stream that is not currently under capture): => If current thread is capturing in relaxed mode, calls are allowed => If any stream in current/concurrent thread is capturing in global mode, calls are not allowed => If any stream in current thread is capturing in ThreadLocal mode, calls are not allowed For Stream that is currently under capture => calls are not allowed => Any call that is not allowed during capture invalidates the capture sequence => It is invalid to call synchronous APIs during capture. Synchronous APIs, such as hipMemcpy(), enqueue work to the legacy stream and synchronize it before returning. Change-Id: I201c6e63e1a5d93fd416a3b520264c0fdbe31237 [ROCm/clr commit: 19b4660cbb84f3c5b61a744700678f5bd0e1425a] --- .../clr/hipamd/src/hip_device_runtime.cpp | 1 + projects/clr/hipamd/src/hip_event.cpp | 10 +++-- projects/clr/hipamd/src/hip_graph.cpp | 12 +++--- projects/clr/hipamd/src/hip_internal.hpp | 34 ++++++++++++++-- projects/clr/hipamd/src/hip_stream.cpp | 39 ++++++++++++------- 5 files changed, 72 insertions(+), 24 deletions(-) diff --git a/projects/clr/hipamd/src/hip_device_runtime.cpp b/projects/clr/hipamd/src/hip_device_runtime.cpp index b199bcdde8..32fc26fbd4 100644 --- a/projects/clr/hipamd/src/hip_device_runtime.cpp +++ b/projects/clr/hipamd/src/hip_device_runtime.cpp @@ -609,6 +609,7 @@ hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config) { hipError_t hipDeviceSynchronize() { HIP_INIT_API(hipDeviceSynchronize); + CHECK_SUPPORTED_DURING_CAPTURE(); constexpr bool kDoWaitForCpu = true; hip::getCurrentDevice()->SyncAllStreams(kDoWaitForCpu); HIP_RETURN(hipSuccess); diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index f527766e4b..8a49bfd8e9 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -434,9 +434,8 @@ hipError_t hipEventSynchronize(hipEvent_t event) { hip::Event* e = reinterpret_cast(event); hip::Stream* s = reinterpret_cast(e->GetCaptureStream()); if ((s != nullptr) && (s->GetCaptureStatus() == hipStreamCaptureStatusActive)) { - if (s->IsEventCaptured(event) == false) { - return HIP_RETURN(hipErrorStreamCaptureUnsupported); - } + s->SetCaptureStatus(hipStreamCaptureStatusInvalidated); + return HIP_RETURN(hipErrorCapturedEvent); } if (hip::Stream::StreamCaptureOngoing(e->GetCaptureStream()) == true) { HIP_RETURN(hipErrorStreamCaptureUnsupported); @@ -455,6 +454,11 @@ hipError_t ihipEventQuery(hipEvent_t event) { } hip::Event* e = reinterpret_cast(event); + hip::Stream* s = reinterpret_cast(e->GetCaptureStream()); + if ((s != nullptr) && (s->GetCaptureStatus() == hipStreamCaptureStatusActive)) { + s->SetCaptureStatus(hipStreamCaptureStatusInvalidated); + return HIP_RETURN(hipErrorCapturedEvent); + } return e->query(); } diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 30ac1fc33f..7481b3b7c0 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -961,7 +961,7 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) { if (s->GetCaptureStatus() == hipStreamCaptureStatusNone) { return hipErrorIllegalState; } - // Capture must be ended on the same stream in which it was initiated + // Capture must be ended on the same stream in which it was initiated if (!s->IsOriginStream()) { return hipErrorStreamCaptureUnmatched; } @@ -978,15 +978,17 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) { amd::ScopedLock lock(g_captureStreamsLock); g_captureStreams.erase(std::find(g_captureStreams.begin(), g_captureStreams.end(), s)); } + { + amd::ScopedLock lock(g_streamSetLock); + g_allCapturingStreams.erase( + std::find(g_allCapturingStreams.begin(), g_allCapturingStreams.end(), s)); + } // If capture was invalidated, due to a violation of the rules of stream capture if (s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) { *pGraph = nullptr; return hipErrorStreamCaptureInvalidated; } - { - amd::ScopedLock lock(g_streamSetLock); - g_allCapturingStreams.erase(std::find(g_allCapturingStreams.begin(), g_allCapturingStreams.end(), s)); - } + // check if all parallel streams have joined // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do // not result in hipErrorStreamCaptureUnjoined diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 643d2c234c..d93c445d4d 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -180,23 +180,48 @@ const char* ihipGetErrorName(hipError_t hip_error); } \ } while (0); +// During stream capture some actions, such as a call to hipMalloc, may be unsafe and prohibited +// during capture. It is allowed only in relaxed mode. #define CHECK_STREAM_CAPTURE_SUPPORTED() \ if (hip::tls.stream_capture_mode_ == hipStreamCaptureModeThreadLocal) { \ if (hip::tls.capture_streams_.size() != 0) { \ + for (auto stream : hip::tls.capture_streams_) { \ + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); \ + } \ HIP_RETURN(hipErrorStreamCaptureUnsupported); \ } \ } else if (hip::tls.stream_capture_mode_ == hipStreamCaptureModeGlobal) { \ if (hip::tls.capture_streams_.size() != 0) { \ + for (auto stream : hip::tls.capture_streams_) { \ + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); \ + } \ HIP_RETURN(hipErrorStreamCaptureUnsupported); \ } \ if (g_captureStreams.size() != 0) { \ + for (auto stream : g_captureStreams) { \ + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); \ + } \ HIP_RETURN(hipErrorStreamCaptureUnsupported); \ } \ } -// Sync APIs cannot be called when stream capture is active +// Device sync is not supported during capture +#define CHECK_SUPPORTED_DURING_CAPTURE() \ + if (!g_allCapturingStreams.empty()) { \ + for (auto stream : g_allCapturingStreams) { \ + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); \ + } \ + return hipErrorStreamCaptureUnsupported; \ + } + +// Sync APIs like hipMemset, hipMemcpy etc.. cannot be called when stream capture is active +// for all capture modes hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal and +// hipStreamCaptureModeRelaxed #define CHECK_STREAM_CAPTURING() \ - if (!g_captureStreams.empty()) { \ + if (!g_allCapturingStreams.empty()) { \ + for (auto stream : g_allCapturingStreams) { \ + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); \ + } \ return hipErrorStreamCaptureImplicit; \ } @@ -207,6 +232,10 @@ const char* ihipGetErrorName(hipError_t hip_error); hipStreamCaptureStatusActive) { \ hipError_t status = hip::capture##name(stream, ##__VA_ARGS__); \ return status; \ + } else if (stream != nullptr && \ + reinterpret_cast(stream)->GetCaptureStatus() == \ + hipStreamCaptureStatusInvalidated) { \ + return hipErrorStreamCaptureInvalidated; \ } #define PER_THREAD_DEFAULT_STREAM(stream) \ @@ -310,7 +339,6 @@ public: static bool StreamCaptureBlocking(); static void Destroy(hip::Stream* stream); - /// Check Stream Capture status to make sure it is done static bool StreamCaptureOngoing(hipStream_t hStream); diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index a78fc093d6..c8a35a53ab 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -130,21 +130,34 @@ bool Stream::StreamCaptureBlocking() { bool Stream::StreamCaptureOngoing(hipStream_t hStream) { hip::Stream* s = reinterpret_cast(hStream); - // Allow capture to be less restrictive one one changes the stream capture interaction - // mode for the thread - if (hip::tls.stream_capture_mode_ == hipStreamCaptureModeRelaxed) { - return false; - } - // 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) { + if (s != nullptr && s->GetCaptureStatus() == hipStreamCaptureStatusNone) { + // If current thread is capturing in relaxed mode + if (hip::tls.stream_capture_mode_ == hipStreamCaptureModeRelaxed) { + return false; + } + // If any stream in current/concurrent thread is capturing in global mode amd::ScopedLock lock(g_captureStreamsLock); - return (g_captureStreams.empty() == true && hip::tls.capture_streams_.empty()) ? false : true; - } - else { - amd::ScopedLock lock(g_streamSetLock); - return (g_allCapturingStreams.find(s) == g_allCapturingStreams.end() ? false : true); + if (!g_captureStreams.empty()) { + for (auto stream : hip::g_captureStreams) { + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); + } + return true; + } + // If any stream in current thread is capturing in ThreadLocal mode + if (!hip::tls.capture_streams_.empty()) { + for (auto stream : hip::tls.capture_streams_) { + stream->SetCaptureStatus(hipStreamCaptureStatusInvalidated); + } + return true; + } + return false; + } else if (s != nullptr && s->GetCaptureStatus() == hipStreamCaptureStatusActive) { + s->SetCaptureStatus(hipStreamCaptureStatusInvalidated); + return true; + } else if (s != nullptr && s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) { + return true; } + return false; } // ================================================================================================