SWDEV-360031 - Erase destroyed stream from capture

Parallel streams can be destroyed during capture. Make sure runtime
removes the destoryed streams from the list of parallel captured streams

Change-Id: I46b6cbb951e8711f0bf02f2826c3b890b2541ffe


[ROCm/clr commit: 38ad6234f9]

projects/clr/hipamd/src/hip_event.cpp

@@ -344,6 +344,9 @@ hipError_t hipEventDestroy(hipEvent_t event) {
   }
 
   hip::Event* e = reinterpret_cast<hip::Event*>(event);
+  if (e->GetCaptureStream() != nullptr) {
+    reinterpret_cast<hip::Stream*>(e->GetCaptureStream())->EraseCaptureEvent(event);
+  }
   delete e;
   HIP_RETURN(hipSuccess);
 }

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<hipGraphNode_t> nodesPrevToRecorded_;
 

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<hipGraphNode_t> lastCapturedNodes_;
     /// dependencies removed via API hipStreamUpdateCaptureDependencies
@@ -258,7 +258,7 @@ namespace hip {
     /// Derived streams/Paralell branches from the origin stream
     std::vector<hipStream_t> parallelCaptureStreams_;
     /// Capture events
-    std::vector<hipEvent_t> captureEvents_;
+    std::unordered_set<hipEvent_t> 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

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<hip::Event*>(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<hip::Stream*>(stream);
-
+  if (s->GetCaptureStatus() != hipStreamCaptureStatusNone) {
+    if (s->GetParentStream() != nullptr) {
+      reinterpret_cast<hip::Stream*>(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<CallbackData*>(user_data);
   int offset = data->previous_read_index % IPC_SIGNALS_PER_EVENT;