From 7450910e534e3aa284510f6a8a1e215b262a0840 Mon Sep 17 00:00:00 2001 From: "systems-assistant[bot]" <221163467+systems-assistant[bot]@users.noreply.github.com> Date: Thu, 13 Nov 2025 11:13:40 +0100 Subject: [PATCH] SWDEV-548241 - Add missing destroy calls in graph tests (#520) Co-authored-by: Vladana Stojiljkovic --- .../hip-tests/catch/unit/graph/hipGraph.cc | 2 + .../unit/graph/hipGraphAddChildGraphNode.cc | 6 ++- .../graph/hipGraphChildGraphNodeGetGraph.cc | 1 + .../catch/unit/graph/hipGraphDestroyNode.cc | 2 +- .../catch/unit/graph/hipGraphExecGetFlags.cc | 2 + .../graph/hipGraphExecMemcpyNodeSetParams.cc | 10 ++-- .../hipGraphExecMemcpyNodeSetParams1D.cc | 12 ++--- .../catch/unit/graph/hipGraphExecUpdate.cc | 1 + .../catch/unit/graph/hipGraphGetNodes_old.cc | 1 + .../unit/graph/hipGraphGetRootNodes_old.cc | 2 +- .../unit/graph/hipGraphHostNodeSetParams.cc | 1 + .../catch/unit/graph/hipGraphInstantiate.cc | 1 + .../graph/hipGraphInstantiateWithFlags.cc | 2 + .../graph/hipGraphInstantiateWithParams.cc | 12 ++--- .../unit/graph/hipGraphNodeFindInClone.cc | 6 ++- .../catch/unit/graph/hipGraphNodeGetType.cc | 7 +++ .../catch/unit/graph/hipGraphNodeSetParams.cc | 3 ++ .../graph/hipGraphRemoveDependencies_old.cc | 3 ++ .../catch/unit/graph/hipStreamBeginCapture.cc | 48 ++++++++++++++++++- .../graph/hipStreamBeginCaptureToGraph.cc | 4 +- .../unit/graph/hipStreamBeginCapture_old.cc | 22 +++++++++ .../hipStreamPerThread_Event.cc | 1 + 22 files changed, 120 insertions(+), 29 deletions(-) diff --git a/projects/hip-tests/catch/unit/graph/hipGraph.cc b/projects/hip-tests/catch/unit/graph/hipGraph.cc index 703d46d169..a44e4dad0e 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraph.cc +++ b/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(stop - start); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc index adc6588d6b..75382d073b 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc +++ b/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(A_d, B_d, nullptr, A_h, B_h, nullptr, false); } /* @@ -219,7 +223,6 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_CloneChildGraph") { HipTest::initArrays(&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)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc b/projects/hip-tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc index c6e781bd5a..72444b6263 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc +++ b/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)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc index fb8ef3d2e5..920a38cec0 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc +++ b/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}; diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecGetFlags.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecGetFlags.cc index 6f3af7ccf3..9c6b5b2337 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecGetFlags.cc +++ b/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)); } /** diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams.cc index b8beaa3b37..d98d8d2449 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams.cc +++ b/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, ¶ms), + hipErrorInvalidValue); } - params.kind = set_dir; - HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams(graph_exec, node, ¶ms), hipErrorInvalidValue); HIP_CHECK(hipGraphExecDestroy(graph_exec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipHostFree(host)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc index 3e14d54b1f..2cea42c65c 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc +++ b/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)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc index ffd4e30a60..3616bd36aa 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc +++ b/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(A_d, nullptr, nullptr, A_h, nullptr, nullptr, false); } /** diff --git a/projects/hip-tests/catch/unit/graph/hipGraphGetNodes_old.cc b/projects/hip-tests/catch/unit/graph/hipGraphGetNodes_old.cc index 1053b26ce3..4dd703c501 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphGetNodes_old.cc +++ b/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") { diff --git a/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes_old.cc b/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes_old.cc index 08d933a03b..166e6e1e08 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes_old.cc +++ b/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}; diff --git a/projects/hip-tests/catch/unit/graph/hipGraphHostNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphHostNodeSetParams.cc index 51d505b1eb..e526fce205 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphHostNodeSetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphHostNodeSetParams.cc @@ -160,6 +160,7 @@ TEST_CASE("Unit_hipGraphHostNodeSetParams_ClonedGraphWithHostNode") { HipTest::freeArrays(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)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphInstantiate.cc b/projects/hip-tests/catch/unit/graph/hipGraphInstantiate.cc index 0fdb4b59bb..4d63692f8a 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiate.cc +++ b/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)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc index e44e048dca..f44e712c52 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc +++ b/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)); } } /* diff --git a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc index 4ed448d3dc..1e724dca90 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc +++ b/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)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphNodeFindInClone.cc b/projects/hip-tests/catch/unit/graph/hipGraphNodeFindInClone.cc index 5c8cb8956b..d06a06a83d 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphNodeFindInClone.cc +++ b/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 diff --git a/projects/hip-tests/catch/unit/graph/hipGraphNodeGetType.cc b/projects/hip-tests/catch/unit/graph/hipGraphNodeGetType.cc index a43ec63b3d..0d5e57be98 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphNodeGetType.cc +++ b/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") { diff --git a/projects/hip-tests/catch/unit/graph/hipGraphNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphNodeSetParams.cc index 013f37a416..5226b00d01 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphNodeSetParams.cc +++ b/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; diff --git a/projects/hip-tests/catch/unit/graph/hipGraphRemoveDependencies_old.cc b/projects/hip-tests/catch/unit/graph/hipGraphRemoveDependencies_old.cc index 24308b2557..0728e00095 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphRemoveDependencies_old.cc +++ b/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); } diff --git a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc index 2721da0504..cb70192049 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc +++ b/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, diff --git a/projects/hip-tests/catch/unit/graph/hipStreamBeginCaptureToGraph.cc b/projects/hip-tests/catch/unit/graph/hipStreamBeginCaptureToGraph.cc index 13b6e6fa04..4021584047 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamBeginCaptureToGraph.cc +++ b/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") { diff --git a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture_old.cc b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture_old.cc index 6c6e83e20b..86cf9a59a0 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture_old.cc +++ b/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 */ diff --git a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc b/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc index 61d0a182b6..552725f099 100644 --- a/projects/hip-tests/catch/unit/streamperthread/hipStreamPerThread_Event.cc +++ b/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)); }