From 8b89aab0d4b4ac26586b9a43b05d4a08e765afc3 Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Thu, 28 Mar 2024 06:14:19 +0000 Subject: [PATCH] SWDEV-436126 - Fix Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread global mode Change-Id: I766ff33485a49df98380cfc8a17bbd9db7548faa [ROCm/hip-tests commit: 3641ca591c781b9497a351c3349ed681a648347a] --- .../hip-tests/catch/unit/graph/hipStreamBeginCapture.cc | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc index 12d82ac09b..a74047aecc 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc +++ b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc @@ -1401,7 +1401,12 @@ static void captureStrmThread(hipGraph_t* graph, int* Ah, int* Ad, int* Bh, int* std::thread t1(strmSyncThread, Ah, Ad, Bh, Bd, BLOCKSIZE, error); t1.join(); myadd<<>>(Ad, Bd); - HIP_CHECK(hipStreamEndCapture(stream[0], graph)); // End Capture + if (flag == hipStreamCaptureModeGlobal) { + HIP_CHECK_ERROR(hipStreamEndCapture(stream[0], graph), + hipErrorStreamCaptureInvalidated); // End Capture + } else { + HIP_CHECK(hipStreamEndCapture(stream[0], graph)); // End Capture + } } TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") { @@ -1445,13 +1450,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") { REQUIRE(error == hipErrorStreamCaptureUnsupported); } #endif -#if HT_AMD SECTION("Capture Flag = hipStreamCaptureModeGlobal Multithreaded") { captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(), Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE, hipStreamCaptureModeGlobal, &error); REQUIRE(error == hipErrorStreamCaptureUnsupported); } -#endif SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Multithreaded") { captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(), Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE, hipStreamCaptureModeThreadLocal, &error);