diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 99603557b8..7f384eff65 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -344,6 +344,9 @@ hipError_t hipEventDestroy(hipEvent_t event) { } hip::Event* e = reinterpret_cast(event); + if (e->GetCaptureStream() != nullptr) { + reinterpret_cast(e->GetCaptureStream())->EraseCaptureEvent(event); + } delete e; HIP_RETURN(hipSuccess); } diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp index bd8e70be1e..e08ea33f66 100644 --- a/projects/clr/hipamd/src/hip_event.hpp +++ b/projects/clr/hipamd/src/hip_event.hpp @@ -31,7 +31,7 @@ protected: public: StreamCallback(void* userData) : userData_(userData) {} - + virtual void CL_CALLBACK callback() = 0; }; @@ -93,7 +93,7 @@ class Event { /// event recorded on stream where capture is active bool onCapture_; /// capture stream where event is recorded - hipStream_t captureStream_; + hipStream_t captureStream_ = nullptr; /// Previous captured nodes before event record std::vector nodesPrevToRecorded_; diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 9ff0690fc9..603390f96c 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -250,7 +250,7 @@ namespace hip { bool originStream_; /// Origin sream has no parent. Parent stream for the derived captured streams with event /// dependencies - hipStream_t parentStream_; + hipStream_t parentStream_ = nullptr; /// Last graph node captured in the stream std::vector lastCapturedNodes_; /// dependencies removed via API hipStreamUpdateCaptureDependencies @@ -258,7 +258,7 @@ namespace hip { /// Derived streams/Paralell branches from the origin stream std::vector parallelCaptureStreams_; /// Capture events - std::vector captureEvents_; + std::unordered_set captureEvents_; unsigned long long captureID_; public: Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false, @@ -361,8 +361,20 @@ namespace hip { } /// Get Capture ID unsigned long long GetCaptureID() { return captureID_; } - void SetCaptureEvent(hipEvent_t e) { captureEvents_.push_back(e); } + void SetCaptureEvent(hipEvent_t e) { captureEvents_.emplace(e); } + void EraseCaptureEvent(hipEvent_t e) { + auto it = captureEvents_.find(e); + if (it != captureEvents_.end()) { + captureEvents_.erase(it); + } + } void SetParallelCaptureStream(hipStream_t s) { parallelCaptureStreams_.push_back(s); } + void EraseParallelCaptureStream(hipStream_t s) { + auto it = std::find(parallelCaptureStreams_.begin(), parallelCaptureStreams_.end(), s); + if (it != parallelCaptureStreams_.end()) { + parallelCaptureStreams_.erase(it); + } + } }; /// HIP Device class diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index d65f738dda..9f6a12bf3a 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -53,6 +53,7 @@ Stream::~Stream() { } } +// ================================================================================================ hipError_t Stream::EndCapture() { for (auto event : captureEvents_) { hip::Event* e = reinterpret_cast(event); @@ -73,6 +74,7 @@ hipError_t Stream::EndCapture() { return hipSuccess; } + // ================================================================================================ bool Stream::Create() { amd::CommandQueue::Priority p; @@ -468,7 +470,12 @@ hipError_t hipStreamDestroy(hipStream_t stream) { HIP_RETURN(hipErrorContextIsDestroyed); } hip::Stream* s = reinterpret_cast(stream); - + if (s->GetCaptureStatus() != hipStreamCaptureStatusNone) { + if (s->GetParentStream() != nullptr) { + reinterpret_cast(s->GetParentStream())->EraseParallelCaptureStream(stream); + } + auto error = s->EndCapture(); + } s->GetDevice()->RemoveStreamFromPools(s); amd::ScopedLock lock(g_captureStreamsLock); @@ -486,6 +493,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) { HIP_RETURN(hipSuccess); } +// ================================================================================================ void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_data) { CallbackData* data = reinterpret_cast(user_data); int offset = data->previous_read_index % IPC_SIGNALS_PER_EVENT;