SWDEV-548241 - Add missing destroy calls in graph tests (#520)

Co-authored-by: Vladana Stojiljkovic <Vladana.Stojiljkovic@amd.com>

projects/hip-tests/catch/unit/graph/hipGraph.cc

@@ -95,6 +95,8 @@ static void hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outpu
                        outputVec_d, result_d);
     HIP_CHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double), hipMemcpyDefault, stream1));
     HIP_CHECK(hipStreamSynchronize(stream1));
+    HIP_CHECK(hipStreamSynchronize(stream2));
+    HIP_CHECK(hipStreamSynchronize(stream3));
   }
   auto stop = std::chrono::high_resolution_clock::now();
   auto result = std::chrono::duration<double, std::milli>(stop - start);

projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc

@@ -109,6 +109,10 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_Negative") {
     REQUIRE(hipGraphAddChildGraphNode(&childGraphNode1, graph, nullptr, 10, childgraph1) ==
             hipErrorInvalidValue);
   }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(childgraph1));
+  HipTest::freeArrays<int>(A_d, B_d, nullptr, A_h, B_h, nullptr, false);
 }
 
 /*
@@ -219,7 +223,6 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_CloneChildGraph") {
   HipTest::initArrays<int>(&A_d, &B_d, nullptr, &A_h, &B_h, nullptr, N, false);
 
   HIP_CHECK(hipGraphCreate(&graph, 0));
-  HIP_CHECK(hipGraphCreate(&clonedgraph, 0));
   hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1;
   hipStream_t streamForGraph;
   HIP_CHECK(hipStreamCreate(&streamForGraph));
@@ -253,6 +256,7 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_CloneChildGraph") {
   HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(childgraph1));
   HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
   HIP_CHECK(hipStreamDestroy(streamForGraph));
 }
 

projects/hip-tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc

@@ -144,5 +144,6 @@ TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Negative") {
   }
 #endif
   HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipGraphDestroy(childgraph1));
 }

projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc

@@ -83,6 +83,7 @@ TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") {
   HIP_CHECK(hipGraphCreate(&graph, 0));
   HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
   REQUIRE(hipGraphDestroyNode(memsetNode) == hipSuccess);
+  HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipFree(pOutBuff_d));
 }
 
@@ -201,7 +202,6 @@ TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph") {
   hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
   hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
   HIP_CHECK(hipGraphCreate(&graph, 0));
-  HIP_CHECK(hipGraphCreate(&clonedgraph, 0));
   // Create graph with no dependencies
   for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
     void* kernelArgs[] = {nullptr};

projects/hip-tests/catch/unit/graph/hipGraphExecGetFlags.cc

@@ -60,6 +60,7 @@ TEST_CASE("Unit_hipGraphExecGetFlags_Negative") {
   HIP_CHECK(
       hipGraphInstantiateWithFlags(&graphExec, graph, hipGraphInstantiateFlagAutoFreeOnLaunch));
   HIP_CHECK_ERROR(hipGraphExecGetFlags(nullptr, &flags), hipErrorInvalidValue);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(graph));
 }
 
@@ -136,6 +137,7 @@ TEST_CASE("Unit_hipGraphExecGetFlags_positive") {
     REQUIRE(flags == hipGraphInstantiateFlagUseNodePriority);
   }
 #endif
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(graph));
 }
 /**

projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams.cc

@@ -256,14 +256,12 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams_Negative_Changing_Memcpy_Directi
 
   const auto set_dir = GENERATE(hipMemcpyHostToHost, hipMemcpyHostToDevice, hipMemcpyDeviceToHost,
                                 hipMemcpyDeviceToDevice, hipMemcpyDefault);
-  if (dir == set_dir) {
-    HIP_CHECK(hipGraphExecDestroy(graph_exec));
-    HIP_CHECK(hipGraphDestroy(graph));
-    return;
+  if (dir != set_dir) {
+    params.kind = set_dir;
+    HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams(graph_exec, node, &params),
+                    hipErrorInvalidValue);
   }
 
-  params.kind = set_dir;
-  HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams(graph_exec, node, &params), hipErrorInvalidValue);
   HIP_CHECK(hipGraphExecDestroy(graph_exec));
   HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipHostFree(host));

projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc

@@ -240,16 +240,12 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Negative_Changing_Memcpy_Direc
 
   const auto set_dir = GENERATE(hipMemcpyHostToHost, hipMemcpyHostToDevice, hipMemcpyDeviceToHost,
                                 hipMemcpyDeviceToDevice, hipMemcpyDefault);
-  if (dir == set_dir) {
-    HIP_CHECK(hipGraphExecDestroy(graph_exec));
-    HIP_CHECK(hipGraphDestroy(graph));
-    return;
+  if (dir != set_dir) {
+    HIP_CHECK_ERROR(
+        hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dst, src, sizeof(int), set_dir),
+        hipErrorInvalidValue);
   }
 
-  HIP_CHECK_ERROR(
-      hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dst, src, sizeof(int), set_dir),
-      hipErrorInvalidValue);
-
   HIP_CHECK(hipGraphExecDestroy(graph_exec));
   HIP_CHECK(hipGraphDestroy(graph));
 

projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc

@@ -119,6 +119,7 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") {
   HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipGraphDestroy(graph2));
   HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HipTest::freeArrays<int>(A_d, nullptr, nullptr, A_h, nullptr, nullptr, false);
 }
 
 /**

projects/hip-tests/catch/unit/graph/hipGraphGetNodes_old.cc

@@ -296,6 +296,7 @@ TEST_CASE("Unit_hipGraphGetNodes_ParamValidation") {
     HIP_CHECK(hipGraphCreate(&emptyGraph, 0));
     HIP_CHECK(hipGraphGetNodes(emptyGraph, nullptr, &numNodes));
     REQUIRE(numNodes == 0);
+    HIP_CHECK(hipGraphDestroy(emptyGraph));
   }
 
   SECTION("numNodes less than actual number of nodes") {

projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes_old.cc

@@ -321,6 +321,7 @@ TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") {
     HIP_CHECK(hipGraphCreate(&emptyGraph, 0));
     HIP_CHECK(hipGraphGetRootNodes(emptyGraph, nullptr, &numRootNodes));
     REQUIRE(numRootNodes == 0);
+    HIP_CHECK(hipGraphDestroy(emptyGraph));
   }
 
   SECTION("numRootNodes less than actual number of nodes") {
@@ -394,7 +395,6 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes_ClonedGrph") {
   hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES];
   hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES];
   HIP_CHECK(hipGraphCreate(&graph, 0));
-  HIP_CHECK(hipGraphCreate(&clonedgraph, 0));
   // Create graph with no dependencies
   for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) {
     void* kernelArgs[] = {nullptr};

projects/hip-tests/catch/unit/graph/hipGraphHostNodeSetParams.cc

@@ -160,6 +160,7 @@ TEST_CASE("Unit_hipGraphHostNodeSetParams_ClonedGraphWithHostNode") {
   HipTest::freeArrays<int>(A_d, nullptr, C_d, A_h, nullptr, C_h, false);
   HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
   HIP_CHECK(hipStreamDestroy(streamForGraph));
 }
 

projects/hip-tests/catch/unit/graph/hipGraphInstantiate.cc

@@ -75,6 +75,7 @@ TEST_CASE("Unit_hipGraphInstantiate_Negative") {
   SECTION("Pass pGraphExec as un-initialize") {
     ret = hipGraphInstantiate(&gExec, graph, nullptr, nullptr, 0);
     REQUIRE(hipSuccess == ret);
+    HIP_CHECK(hipGraphExecDestroy(gExec));
   }
   HIP_CHECK(hipGraphDestroy(graph));
 }

projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc

@@ -70,6 +70,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") {
     hipGraph_t graph;
     HIP_CHECK(hipGraphCreate(&graph, 0));
     REQUIRE(hipGraphInstantiateWithFlags(nullptr, graph, 0) == hipErrorInvalidValue);
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 
   SECTION("Passing nullptr to graph") {
@@ -82,6 +83,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") {
     HIP_CHECK(hipGraphCreate(&graph, 0));
     hipGraphExec_t graphExec;
     REQUIRE(hipGraphInstantiateWithFlags(&graphExec, graph, 10) != hipSuccess);
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 }
 /*

projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc

@@ -197,17 +197,15 @@ void GraphInstantiateWithParams_StreamCapture() {
   HIP_CHECK(hipMalloc(&C_d, Nbytes));
   REQUIRE(A_d != nullptr);
   REQUIRE(C_d != nullptr);
-  HIP_CHECK(hipGraphCreate(&graph, 0));
-
 
   HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
+
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
+  HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
+
   constexpr unsigned blocks = 512;
   constexpr unsigned threadsPerBlock = 256;
-
-  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
-  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
-
-  HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
   hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d,
                      C_d, N);
   HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));

projects/hip-tests/catch/unit/graph/hipGraphNodeFindInClone.cc

@@ -183,8 +183,6 @@ void hipGraphNodeFindInClone_Func(bool ModifyOrigGraph = false) {
   HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
 
   if (ModifyOrigGraph) {
-    // Cloned the graph
-    HIP_CHECK(hipGraphClone(&clonedgraph, graph));
     // Modify Original graph by adding new dependency
     HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0, C_d, B_d, Nbytes,
                                       hipMemcpyDeviceToHost));
@@ -259,6 +257,10 @@ void hipGraphNodeFindInClone_DoubleClone(bool ModifyOrigGraph = false) {
   hipGraphNode_t clonedgraphnode;
   REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, memcpyH2D_A, clonedgraph_1) ==
           hipErrorInvalidValue);
+  HIP_CHECK(hipGraphDestroy(clonedgraph_1));
+  HIP_CHECK(hipGraphDestroy(clonedgraph));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
 }
 /**
  * Test Description

projects/hip-tests/catch/unit/graph/hipGraphNodeGetType.cc

@@ -74,6 +74,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_Negative") {
     HIP_CHECK(hipGraphCreate(&graph, 0));
     HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr, 0));
     REQUIRE(hipGraphNodeGetType(memcpyNode, nullptr) == hipErrorInvalidValue);
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 
   SECTION("Pass invalid node") {
@@ -114,8 +115,11 @@ TEST_CASE("Unit_hipGraphNodeGetType_Functional") {
     HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType));
     REQUIRE(nodeType == hipGraphNodeTypeEmpty);
   }
+
+  HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipStreamDestroy(stream));
   HIP_CHECK(hipEventDestroy(event));
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
 }
 /**
  * Functional Test for hipGraphNodeGetType API
@@ -397,6 +401,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfClonedGraph_NodeTypeInThread") {
   SECTION("Cloned Graph Node Type") {
     HIP_CHECK(hipGraphClone(&clonedGraph, graph));
     ChkNodeType(clonedGraph, &numNode);
+    HIP_CHECK(hipGraphDestroy(clonedGraph));
   }
   // Thread
   SECTION("Node Type In The Thread") {
@@ -515,6 +520,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_NodeTypeOfChildGraph") {
   HIP_CHECK(hipGraphChildGraphNodeGetGraph(childGraphNode, &getGraph));
   ChkNodeType(getGraph, &numNodeChild);
 
+  HIP_CHECK(hipStreamSynchronize(stream2));
   HIP_CHECK(hipStreamDestroy(stream1));
   HIP_CHECK(hipEventDestroy(event1));
   HIP_CHECK(hipStreamDestroy(stream2));
@@ -638,6 +644,7 @@ TEST_CASE("Unit_hipGraphNodeGetType_ClonedGraph_InThread_WithDependencies") {
   SECTION("Cloned Graph Node Type") {
     HIP_CHECK(hipGraphClone(&clonedGraph, graph));
     ChkNodeTypeWithDependency(clonedGraph, Parent);
+    HIP_CHECK(hipGraphDestroy(clonedGraph));
   }
   // Thread
   SECTION("Node Type In The Thread") {

projects/hip-tests/catch/unit/graph/hipGraphNodeSetParams.cc

@@ -88,6 +88,8 @@ TEST_CASE("Unit_hipGraphNodeSetParams_Negative_Parameters") {
 
   HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
 }
 
 /**
@@ -130,6 +132,7 @@ TEST_CASE("Unit_hipGraphNodeSetParams_Positive") {
   for (int i = 0; i < N; i++) {
     REQUIRE(A_h[i] == 99);
   }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
 
   hipGraphNodeParams node_params2 = {};
   node_params2.type = hipGraphNodeTypeMemset;

projects/hip-tests/catch/unit/graph/hipGraphRemoveDependencies_old.cc

@@ -256,6 +256,9 @@ TEST_CASE("Unit_hipGraphRemoveDependencies_Func_StrmCapture") {
   HIP_CHECK(hipStreamDestroy(stream1));
   HIP_CHECK(hipStreamDestroy(stream2));
   HIP_CHECK(hipStreamDestroy(stream3));
+  HIP_CHECK(hipEventDestroy(forkStreamEvent));
+  HIP_CHECK(hipEventDestroy(memsetEvent1));
+  HIP_CHECK(hipEventDestroy(memsetEvent2));
   HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
 }
 

projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc

@@ -195,6 +195,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") {
     HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
     HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal),
                     hipErrorIllegalState);
+    hipGraph_t graph;
+    HIP_CHECK(hipStreamEndCapture(stream, &graph));
+    HIP_CHECK(hipGraphDestroy(graph));
   }
   SECTION("Creating hipStream with invalid mode") {
     HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)), hipErrorInvalidValue);
@@ -315,6 +318,7 @@ static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t&
   HIP_CHECK(hipGraphExecDestroy(graphExec1));
   HIP_CHECK(hipGraphDestroy(graph2));
   HIP_CHECK(hipGraphDestroy(graph1));
+  HIP_CHECK(hipEventDestroy(event));
 }
 
 /* Local function for colligated stream capture functionality
@@ -653,6 +657,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Multiplestrms") {
     REQUIRE(numNodes1 == 1);
     REQUIRE(numNodes2 == 1);
     REQUIRE(numNodes3 == 1);
+    HIP_CHECK(hipEventDestroy(event));
   }
 
   for (int i = 0; i < 3; i++) {
@@ -754,6 +759,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_DetectingInvalidCapture") {
   HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal),
                   hipErrorIllegalState);
   HIP_CHECK(hipStreamEndCapture(streams[0], &graph));
+  HIP_CHECK(hipGraphDestroy(graph));
 }
 
 /**
@@ -846,6 +852,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_CheckingSyncDuringCapture") {
     HIP_CHECK(hipEventRecord(e, stream));
     HIP_CHECK_ERROR(hipEventQuery(e), hipErrorCapturedEvent);
   }
+
+  hipGraph_t graph;
+  HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph), hipErrorStreamCaptureInvalidated);
 }
 
 /**
@@ -880,21 +889,40 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Concurrent_CheckingSyncDuringCapt
   HIP_CHECK(hipStreamBeginCapture(stream, captureMode));
   SECTION("Synchronize stream during capture") {
     streamSync func;
+    hipGraph_t gr;
     hipError_t expected = hipSuccess;
-    if (captureMode == hipStreamCaptureModeGlobal) expected = hipErrorStreamCaptureUnsupported;
+    hipError_t capture_err = hipSuccess;
+    if (captureMode == hipStreamCaptureModeGlobal) {
+      expected = hipErrorStreamCaptureUnsupported;
+      capture_err = hipErrorStreamCaptureInvalidated;
+    }
 
     std::thread t(std::ref(func), concurrent_stream);
     t.join();
     REQUIRE(func.result_status == expected);
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream, &gr), capture_err);
+    if (capture_err == hipSuccess) {
+      HIP_CHECK(hipGraphDestroy(gr));
+    }
   }
   SECTION("Query stream during capture") {
     streamQuery func;
+    hipGraph_t gr;
     hipError_t expected = hipSuccess;
-    if (captureMode == hipStreamCaptureModeGlobal) expected = hipErrorStreamCaptureUnsupported;
+    hipError_t capture_err = hipSuccess;
+    if (captureMode == hipStreamCaptureModeGlobal) {
+      expected = hipErrorStreamCaptureUnsupported;
+      capture_err = hipErrorStreamCaptureInvalidated;
+    }
 
     std::thread t(std::ref(func), concurrent_stream);
     t.join();
     REQUIRE(func.result_status == expected);
+
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream, &gr), capture_err);
+    if (capture_err == hipSuccess) {
+      HIP_CHECK(hipGraphDestroy(gr));
+    }
   }
   SECTION("Synchronize device during capture") {
     deviceSync func;
@@ -903,6 +931,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Concurrent_CheckingSyncDuringCapt
     std::thread t(std::ref(func));
     t.join();
     REQUIRE(func.result_status == expected);
+    hipGraph_t gr;
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream, &gr), hipErrorStreamCaptureInvalidated);
   }
   SECTION("Synchronize event during capture") {
     eventSync func;
@@ -911,6 +941,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Concurrent_CheckingSyncDuringCapt
     std::thread t(std::ref(func), e);
     t.join();
     REQUIRE(func.result_status == expected);
+    hipGraph_t gr;
+    HIP_CHECK(hipStreamEndCapture(stream, &gr));
+    HIP_CHECK(hipGraphDestroy(gr));
   }
   SECTION("Query for an event during capture") {
     eventQuery func;
@@ -919,6 +952,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Concurrent_CheckingSyncDuringCapt
     std::thread t(std::ref(func), e);
     t.join();
     REQUIRE(func.result_status == expected);
+    hipGraph_t gr;
+    HIP_CHECK(hipStreamEndCapture(stream, &gr));
+    HIP_CHECK(hipGraphDestroy(gr));
   }
 }
 
@@ -961,6 +997,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") {
   SECTION("hipMemset during capture") {
     HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)), hipErrorStreamCaptureImplicit);
   }
+
+  hipGraph_t graph;
+  HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph), hipErrorStreamCaptureInvalidated);
 }
 
 /**
@@ -996,6 +1035,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapwhenCapInProg") {
     HIP_CHECK(hipStreamWaitEvent(stream2, e, 0));
     dummyKernel<<<1, 1, 0, stream2>>>();
     HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined);
+    HIP_CHECK(hipEventDestroy(e));
   }
   SECTION("End strm capture when forked strm still has operations") {
     EventsGuard events_guard(2);
@@ -1524,6 +1564,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") {
                              stream[1]));
     error = hipStreamSynchronize(stream[1]);
     REQUIRE(error == hipErrorStreamCaptureUnsupported);
+    hipGraph_t graph;
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream[0], &graph), hipErrorStreamCaptureInvalidated);
   }
   SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Single Threaded") {
     StreamsGuard stream(2);
@@ -1535,6 +1577,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") {
                              stream[1]));
     error = hipStreamSynchronize(stream[1]);
     REQUIRE(error == hipErrorStreamCaptureUnsupported);
+    hipGraph_t graph;
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream[0], &graph), hipErrorStreamCaptureInvalidated);
   }
   SECTION("Capture Flag = hipStreamCaptureModeGlobal Multithreaded") {
     captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(), Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,

projects/hip-tests/catch/unit/graph/hipStreamBeginCaptureToGraph.cc

@@ -813,6 +813,7 @@ TEST_CASE("Unit_hipStreamBeginCaptureToGraph_ModifyChildpGraph") {
   hipGraphNode_t childGraphNode;
   HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, dependncy.data(), dependncy.size(),
                                       graphChild));
+  HIP_CHECK(hipGraphDestroy(graphChild));
   HIP_CHECK(hipGraphChildGraphNodeGetGraph(childGraphNode, &graphChild));
   // Capture stream into graph
   // Capture Nodes from multiple streams
@@ -907,7 +908,6 @@ TEST_CASE("Unit_hipStreamBeginCaptureToGraph_StateTesting") {
   HIP_CHECK(hipStreamCreate(&stream2));
   HIP_CHECK(hipEventCreate(&e));
   hipStreamCaptureStatus captureStatus = hipStreamCaptureStatusNone;
-  HIP_CHECK(hipGraphCreate(&graph, 0));
   HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus));
   REQUIRE(captureStatus == hipStreamCaptureStatusNone);
   HIP_CHECK(hipStreamBeginCaptureToGraph(stream1, graph, nullptr, nullptr, 0,
@@ -1009,6 +1009,7 @@ TEST_CASE("Unit_hipStreamBeginCaptureToGraph_EndingWhileCaptureInProgress") {
     HIP_CHECK(hipMemcpyAsync(A_d, A_h.data(), Nbytes, hipMemcpyHostToDevice, stream1));
     REQUIRE(hipSuccess == hipStreamEndCapture(stream1, &graph));
     HIP_CHECK(hipEventDestroy(e));
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 
   SECTION("End strm capture when forked strm still has operations") {
@@ -1132,6 +1133,7 @@ static void threadCaptureStart(hipStream_t* streamCapt, hipStream_t* streamFork,
   HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, *streamFork));
   HIP_CHECK(hipEventRecord(e, *streamFork));
   HIP_CHECK(hipStreamWaitEvent(*streamCapt, e, 0));
+  HIP_CHECK(hipEventDestroy(e));
 }
 
 TEST_CASE("Unit_hipStreamBeginCaptureToGraph_CapturePartialInThreads") {

projects/hip-tests/catch/unit/graph/hipStreamBeginCapture_old.cc

@@ -289,6 +289,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative") {
     HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
     ret = hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal);
     REQUIRE(hipErrorIllegalState == ret);
+    hipGraph_t graph;
+    HIP_CHECK(hipStreamEndCapture(stream, &graph));
+    HIP_CHECK(hipGraphDestroy(graph));
   }
   SECTION("Creating hipStream with invalid mode") {
     ret = hipStreamBeginCapture(stream, hipStreamCaptureMode(-1));
@@ -309,6 +312,15 @@ TEST_CASE("Unit_hipStreamBeginCapture_Basic") {
   HIP_CHECK(hipStreamCreate(&s3));
   HIP_CHECK(hipStreamBeginCapture(s3, hipStreamCaptureModeRelaxed));
 
+  hipGraph_t g1, g2, g3;
+  HIP_CHECK(hipStreamEndCapture(s1, &g1));
+  HIP_CHECK(hipStreamEndCapture(s2, &g2));
+  HIP_CHECK(hipStreamEndCapture(s3, &g3));
+
+  HIP_CHECK(hipGraphDestroy(g1));
+  HIP_CHECK(hipGraphDestroy(g2));
+  HIP_CHECK(hipGraphDestroy(g3));
+
   HIP_CHECK(hipStreamDestroy(s1));
   HIP_CHECK(hipStreamDestroy(s2));
   HIP_CHECK(hipStreamDestroy(s3));
@@ -646,6 +658,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_multiplestrms") {
     REQUIRE(numNodes3 == 1);
     HIP_CHECK(hipEventDestroy(event2));
     HIP_CHECK(hipEventDestroy(event1));
+    HIP_CHECK(hipGraphDestroy(graph1));
+    HIP_CHECK(hipGraphDestroy(graph2));
+    HIP_CHECK(hipGraphDestroy(graph3));
   }
   SECTION("Capture Multiple stream with single event") {
     hipEvent_t event1;
@@ -669,6 +684,9 @@ TEST_CASE("Unit_hipStreamBeginCapture_multiplestrms") {
     REQUIRE(numNodes2 == 1);
     REQUIRE(numNodes3 == 1);
     HIP_CHECK(hipEventDestroy(event1));
+    HIP_CHECK(hipGraphDestroy(graph1));
+    HIP_CHECK(hipGraphDestroy(graph2));
+    HIP_CHECK(hipGraphDestroy(graph3));
   }
   HIP_CHECK(hipStreamDestroy(stream3));
   HIP_CHECK(hipStreamDestroy(stream2));
@@ -774,8 +792,12 @@ TEST_CASE("Unit_hipStreamBeginCapture_DetectingInvalidCapture") {
   // Since stream2 is already in capture mode due to event wait
   // hipStreamBeginCapture on stream2 is expected to return error.
   REQUIRE(hipSuccess != hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal));
+  hipGraph_t graph;
+  HIP_CHECK(hipStreamEndCapture(stream1, &graph));
+  HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipStreamDestroy(stream2));
   HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipEventDestroy(event));
 }
 /* Test scenario 12
  */

projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc

@@ -23,6 +23,7 @@ TEST_CASE("Unit_hipStreamPerThread_EventRecord") {
   hipEvent_t event;
   HIP_CHECK(hipEventCreate(&event));
   HIP_CHECK(hipEventRecord(event, hipStreamPerThread));
+  HIP_CHECK(hipEventSynchronize(event));
   HIP_CHECK(hipEventDestroy(event));
 }