AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-443567 - SWDEV-436126 - Fix Prohibited and Unhandled Operations during capture

Browse Source 
=> 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: 19b4660cbb]
This commit is contained in:
Anusha GodavarthySurya
2024-01-25 10:32:08 +00:00
committed by Anusha Godavarthy Surya
parent 4f123e9c10
commit ae296c8fad
5 changed files with 72 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_device_runtime.cpp
+1
View File
@@ -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);
projects/clr/hipamd/src/hip_event.cpp
+7 -3
View File
@@ -434,9 +434,8 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
hip::Event* e = reinterpret_cast<hip::Event*>(event);
hip::Stream* s = reinterpret_cast<hip::Stream*>(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<hip::Event*>(event);
hip::Stream* s = reinterpret_cast<hip::Stream*>(e->GetCaptureStream());
if ((s != nullptr) && (s->GetCaptureStatus() == hipStreamCaptureStatusActive)) {
s->SetCaptureStatus(hipStreamCaptureStatusInvalidated);
return HIP_RETURN(hipErrorCapturedEvent);
}
return e->query();
}
projects/clr/hipamd/src/hip_graph.cpp
+7 -5
View File
@@ -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
projects/clr/hipamd/src/hip_internal.hpp
+31 -3
View File
@@ -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<hip::Stream*>(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);
projects/clr/hipamd/src/hip_stream.cpp
+26 -13
View File
@@ -130,21 +130,34 @@ bool Stream::StreamCaptureBlocking() {
bool Stream::StreamCaptureOngoing(hipStream_t hStream) {
hip::Stream* s = reinterpret_cast<hip::Stream*>(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;
}
// ================================================================================================