From 1af17d765d50bae4ca6ae311fc6354c7374e3900 Mon Sep 17 00:00:00 2001 From: Ajay Date: Mon, 7 Mar 2022 20:35:51 +0000 Subject: [PATCH] SWDEV-315509 - hipStreamEndCapture issues Change-Id: Id95a77ed9ac57f9e904a32947a3d11e3b353be20 --- hipamd/src/hip_graph.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index a4b64f9e1b..5890f6ed79 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -685,9 +685,11 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns s->SetCaptureMode(reinterpret_cast(e->GetCaptureStream())->GetCaptureMode()); s->SetParentStream(e->GetCaptureStream()); s->SetParallelCaptureStream(stream); + } else { + assert(std::find(g_captureStreams.begin(), g_captureStreams.end(), stream) != + g_captureStreams.end() && "capturing stream should be present"); } s->AddCrossCapturedNode(e->GetNodesPrevToRecorded()); - g_captureStreams.push_back(stream); return hipSuccess; } @@ -773,9 +775,9 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { } // If mode is not hipStreamCaptureModeRelaxed, hipStreamEndCapture must be called on the stream // from the same thread + const auto& it = std::find(g_captureStreams.begin(), g_captureStreams.end(), stream); if (s->GetCaptureMode() != hipStreamCaptureModeRelaxed && - std::find(g_captureStreams.begin(), g_captureStreams.end(), stream) == - g_captureStreams.end()) { + it == g_captureStreams.end()) { HIP_RETURN(hipErrorStreamCaptureWrongThread); } // If capture was invalidated, due to a violation of the rules of stream capture @@ -786,7 +788,7 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { // check if all parallel streams have joined // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do // not result in hipErrorStreamCaptureUnjoined - if (s->GetCaptureGraph()->GetLeafNodeCount() != 1) { + if (s->GetCaptureGraph()->GetLeafNodeCount() > 1) { std::vector leafNodes = s->GetCaptureGraph()->GetLeafNodes(); const std::vector& removedDepNodes = s->GetRemovedDependencies(); bool foundInRemovedDep = false; @@ -802,7 +804,8 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { } } *pGraph = s->GetCaptureGraph(); - g_captureStreams.clear(); + // erase the stream and move the elements to the erased spot + g_captureStreams.erase(it); // end capture on all streams/events part of graph capture HIP_RETURN_DURATION(s->EndCapture()); }