diff --git a/catch/include/memcpy3d_tests_common.hh b/catch/include/memcpy3d_tests_common.hh index 58d5be6755..1b8c136e27 100644 --- a/catch/include/memcpy3d_tests_common.hh +++ b/catch/include/memcpy3d_tests_common.hh @@ -181,6 +181,9 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, hipStream_t kernel_stream = null } if constexpr (enable_peer_access) { if (src_device == dst_device) { + if (device_count > 0 && kernel_stream != nullptr && kernel_stream != hipStreamPerThread) { + HIP_CHECK(hipStreamDestroy(kernel_stream)); + } return; } int can_access_peer = 0; @@ -189,6 +192,9 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, hipStream_t kernel_stream = null std::string msg = "Skipped as peer access cannot be enabled between devices " + std::to_string(src_device) + " " + std::to_string(dst_device); HipTest::HIP_SKIP_TEST(msg.c_str()); + if (device_count > 0 && kernel_stream != nullptr && kernel_stream != hipStreamPerThread) { + HIP_CHECK(hipStreamDestroy(kernel_stream)); + } return; } HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0)); @@ -888,4 +894,4 @@ void DrvMemcpy3DArrayDeviceShell(F memcpy_func, const hipStream_t kernel_stream }; PitchedMemoryVerify(host_alloc.ptr(), extent.width, extent.width / sizeof(int), extent.height, extent.depth, f); -} \ No newline at end of file +} diff --git a/catch/include/resource_guards.hh b/catch/include/resource_guards.hh index 91d8b6ffd5..c545fad074 100644 --- a/catch/include/resource_guards.hh +++ b/catch/include/resource_guards.hh @@ -307,8 +307,10 @@ class StreamGuard { break; case Streams::withFlags: HIP_CHECK(hipStreamCreateWithFlags(&stream_, flags_)); + break; case Streams::withPriority: HIP_CHECK(hipStreamCreateWithPriority(&stream_, flags_, priority_)); + break; } } @@ -318,7 +320,7 @@ class StreamGuard { StreamGuard& operator=(StreamGuard&& o) { if (this != &o) { - if (stream_type_ == Streams::created) { + if (stream_type_ >= Streams::created) { static_cast(hipStreamDestroy(stream_)); } @@ -337,7 +339,7 @@ class StreamGuard { } ~StreamGuard() { - if (stream_type_ == Streams::created && stream_ != nullptr) { + if (stream_type_ >= Streams::created && stream_ != nullptr) { static_cast(hipStreamDestroy(stream_)); } } @@ -361,7 +363,9 @@ class EventsGuard { EventsGuard(EventsGuard&&) = delete; ~EventsGuard() { - for (auto& e : events_) static_cast(hipEventDestroy(e)); + for (auto& e : events_) { + static_cast(hipEventDestroy(e)); + } } hipEvent_t& operator[](int index) { return events_[index]; } diff --git a/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc b/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc index a9a66dfe99..87b8f63017 100644 --- a/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc +++ b/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc @@ -374,7 +374,6 @@ static void test_cg_multi_grid_group_type(F kernel_func, int num_devices, int bl HIP_CHECK(hipHostFree(size_host[i])); HIP_CHECK(hipHostFree(thd_rank_host[i])); HIP_CHECK(hipHostFree(is_valid_host[i])); - HIP_CHECK(hipStreamDestroy(stream[i])); } } diff --git a/catch/unit/graph/hipGraph.cc b/catch/unit/graph/hipGraph.cc index 59c552c753..e1c1289584 100644 --- a/catch/unit/graph/hipGraph.cc +++ b/catch/unit/graph/hipGraph.cc @@ -105,6 +105,9 @@ static void hipWithoutGraphs(float* inputVec_h, float* inputVec_d, INFO("Time taken for hipWithoutGraphs : " << std::chrono::duration_cast(result).count() << " millisecs "); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memsetEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); HIP_CHECK(hipStreamDestroy(stream1)); HIP_CHECK(hipStreamDestroy(stream2)); HIP_CHECK(hipStreamDestroy(stream3)); diff --git a/catch/unit/graph/hipGraphClone.cc b/catch/unit/graph/hipGraphClone.cc index 75b83c82a0..14e418c193 100644 --- a/catch/unit/graph/hipGraphClone.cc +++ b/catch/unit/graph/hipGraphClone.cc @@ -208,6 +208,7 @@ void hipGraphClone_Func(bool ModifyOrigGraph = false) { HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); HIP_CHECK(hipStreamSynchronize(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); for (size_t i= 0; i < NElem; i++) { if (C_h[i] != B_h[i]) { diff --git a/catch/unit/graph/hipGraphCloneComplx.cc b/catch/unit/graph/hipGraphCloneComplx.cc index 6f0dd323ef..e8e40af506 100644 --- a/catch/unit/graph/hipGraphCloneComplx.cc +++ b/catch/unit/graph/hipGraphCloneComplx.cc @@ -1302,6 +1302,7 @@ static void hipGraphClone_Test_hipGraphEventRecordNodeSetEvent_and_Exec() { nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(clonedGraphExec, cg.stream)); HIP_CHECK(hipStreamSynchronize(cg.stream)); + HIP_CHECK(hipGraphExecDestroy(clonedGraphExec)); // Verify graph execution result HipTest::checkVectorADD(cg.B_h, cg.C_h, cg.A_h, N); @@ -1332,6 +1333,7 @@ static void hipGraphClone_Test_hipGraphEventRecordNodeSetEvent_and_Exec() { HipTest::checkVectorADD(cg.B_h, cg.C_h, cg.A_h, N); HIP_CHECK(hipEventDestroy(event_end2)); + HIP_CHECK(hipGraphExecDestroy(clonedGraphExec)); } SECTION("Verify hipGraphEventRecordNodeSetEvent & event_end->event_end3") { hipEvent_t event_end3; @@ -1414,9 +1416,9 @@ static void hipGraphClone_Test_hipGraphEventRecordNodeSetEvent_and_Exec() { HipTest::checkVectorADD(cg.B_h, cg.C_h, cg.A_h, N); HIP_CHECK(hipEventDestroy(event_end5)); + HIP_CHECK(hipGraphExecDestroy(clonedGraphExec)); } - HIP_CHECK(hipGraphExecDestroy(clonedGraphExec)); HIP_CHECK(hipGraphDestroy(clonedGraph)); HIP_CHECK(hipGraphDestroy(childgraph)); HIP_CHECK(hipEventDestroy(event_start)); @@ -1518,6 +1520,9 @@ static void hipGraphClone_Test_hipGraphEventWaitNodeSetEvent_and_Exec() { HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_rec_node, event_2)); HIP_CHECK(hipGraphEventWaitNodeSetEvent(event_wait_node, event_2)); + // Destroy clonedGraphExec before instantating a new one + HIP_CHECK(hipGraphExecDestroy(clonedGraphExec)); + // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&clonedGraphExec, clonedGraph, nullptr, nullptr, 0)); @@ -2126,4 +2131,3 @@ TEST_CASE("Unit_hipGraphChild_hipUserObject_hipGraphUserObject") { HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); } - diff --git a/catch/unit/graph/hipGraphCycle.cc b/catch/unit/graph/hipGraphCycle.cc index bb1a7e63ae..f11ea8c43b 100644 --- a/catch/unit/graph/hipGraphCycle.cc +++ b/catch/unit/graph/hipGraphCycle.cc @@ -56,6 +56,7 @@ TEST_CASE("Unit_hipGraph_BasicCyclic1") { HIP_CHECK(hipStreamCreate(&streamForGraph)); REQUIRE(hipErrorInvalidValue == hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); } /** @@ -83,13 +84,15 @@ TEST_CASE("Unit_hipGraph_BasicCyclic2") { HIP_CHECK(hipGraphRemoveDependencies(graph, &emptyNode3, &emptyNode1, 1)); HIP_CHECK(hipStreamCreate(&streamForGraph)); HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } /** * Tests basic functionality of cycle detection in hipGraph APIs by * Adding manual empty nodes - * Cyclic graph, cycle formation first, Remove edge causes disconnected graph which is still + * Cyclic graph, cycle formation first, Remove edge causes disconnected graph which is still * cyclic */ TEST_CASE("Unit_hipGraph_BasicCyclic3") { @@ -115,6 +118,7 @@ TEST_CASE("Unit_hipGraph_BasicCyclic3") { HIP_CHECK(hipStreamCreate(&streamForGraph)); REQUIRE(hipErrorInvalidValue == hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); } /** diff --git a/catch/unit/graph/hipGraphDebugDotPrint.cc b/catch/unit/graph/hipGraphDebugDotPrint.cc index 04e3c874e3..2ebf5ff741 100644 --- a/catch/unit/graph/hipGraphDebugDotPrint.cc +++ b/catch/unit/graph/hipGraphDebugDotPrint.cc @@ -284,6 +284,7 @@ static void hipGraphDebugDotPrint_Functional(const char* fName, free(mem_h); HIP_CHECK(hipFree(mem_d)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipEventDestroy(event)); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipGraphDestroy(childGraph)); diff --git a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol_old.cc b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol_old.cc index b2a92f94db..a2a0fd1d28 100644 --- a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol_old.cc +++ b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol_old.cc @@ -182,6 +182,7 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative") { } HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); } diff --git a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc index 4d3d0189d6..ee05f191b9 100644 --- a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc +++ b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc @@ -181,14 +181,14 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") { SECTION("Changing src allocation device") { if (HipTest::getDeviceCount() < 2) { HipTest::HIP_SKIP_TEST("Test requires two connected GPUs"); - return; + } else { + HIP_CHECK(hipSetDevice(1)); + LinearAllocGuard new_var(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol( + graph_exec, node, SYMBOL(int_device_var), new_var.ptr(), + sizeof(*new_var.ptr()), 0, static_cast(-1)), + hipErrorInvalidValue); } - HIP_CHECK(hipSetDevice(1)); - LinearAllocGuard new_var(LinearAllocs::hipMalloc, sizeof(int)); - HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol( - graph_exec, node, SYMBOL(int_device_var), new_var.ptr(), - sizeof(*new_var.ptr()), 0, static_cast(-1)), - hipErrorInvalidValue); } HIP_CHECK(hipGraphExecDestroy(graph_exec)); diff --git a/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc b/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc index 8533ed4170..123838e80c 100644 --- a/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc +++ b/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc @@ -167,13 +167,13 @@ TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative_Parameters") { SECTION("Changing dst allocation device") { if (HipTest::getDeviceCount() < 2) { HipTest::HIP_SKIP_TEST("Test requires two connected GPUs"); - return; + } else { + HIP_CHECK(hipSetDevice(1)); + LinearAllocGuard new_alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); + params.dst = new_alloc.ptr(); + HIP_CHECK_ERROR(hipGraphExecMemsetNodeSetParams(graph_exec, node, ¶ms), + hipErrorInvalidValue); } - HIP_CHECK(hipSetDevice(1)); - LinearAllocGuard new_alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); - params.dst = new_alloc.ptr(); - HIP_CHECK_ERROR(hipGraphExecMemsetNodeSetParams(graph_exec, node, ¶ms), - hipErrorInvalidValue); } HIP_CHECK(hipGraphExecDestroy(graph_exec)); diff --git a/catch/unit/graph/hipGraphExecUpdate.cc b/catch/unit/graph/hipGraphExecUpdate.cc index bbe00df59a..bbe91ae4b3 100644 --- a/catch/unit/graph/hipGraphExecUpdate.cc +++ b/catch/unit/graph/hipGraphExecUpdate.cc @@ -108,7 +108,6 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") { dependencies.push_back(memsetNode); HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice)); // graphExec was created before memcpyTemp was added to graph. @@ -488,6 +487,7 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_2") { #else REQUIRE(hipSuccess == ret); #endif + HIP_CHECK(hipGraphExecDestroy(graphExec2)); HIP_CHECK(hipGraphDestroy(graph2)); } SECTION("When A node is deleted in GraphExec but not its pair from Graph") { @@ -506,6 +506,7 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_2") { #endif REQUIRE(NULL == hErrorNode_out); + HIP_CHECK(hipGraphExecDestroy(graphExec3)); HIP_CHECK(hipGraphDestroy(graph3)); } HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); diff --git a/catch/unit/graph/hipGraphPerf.cc b/catch/unit/graph/hipGraphPerf.cc index 3f7b2b1b5f..817030f1c0 100644 --- a/catch/unit/graph/hipGraphPerf.cc +++ b/catch/unit/graph/hipGraphPerf.cc @@ -386,6 +386,8 @@ static void checkGraphEventcontinuousKernelCall(const unsigned int kNumNode) { HipTest::checkVectorADD(A_h, B_h, C_h, N); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipEventDestroy(eventstart)); + HIP_CHECK(hipEventDestroy(eventend)); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); @@ -483,6 +485,9 @@ static void checkGraphEventcontinuousKernelCallIn2Blocks( HipTest::checkVectorSUB(A_h, B_h, C_h, N); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipEventDestroy(eventstart)); + HIP_CHECK(hipEventDestroy(eventmid)); + HIP_CHECK(hipEventDestroy(eventend)); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); @@ -582,6 +587,7 @@ static void hipGraph_PerfCheck_hipGraphExecKernelNodeSetParams( HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); // Verify graph execution result HipTest::checkVectorADD(A_h, B_h, C_h, N); @@ -2519,7 +2525,7 @@ static void hipGraph_PerfCheck_hipGraphExecUpdate(const hipStream_t& stream) { * ------------------------ * - Validate hipGraph performance with doorbell set. * - DEBUG_CLR_GRAPH_PACKET_CAPTURE - * 1) Added 2 nodes of MemCpy & a Kernel node and copy back result using memcpy + * 1) Added 2 nodes of MemCpy & a Kernel node and copy back result using memcpy and Instantiate graph & update new graph with similar node structure with api hipGraphExecUpdate and verify the result, the updated node should reflect. i) Check with Multi device case. @@ -2667,7 +2673,7 @@ static void hipGraph_PerfCheck_hipGraphExecUpdate_kernel_inLoop( * ------------------------ * - Validate hipGraph performance with doorbell set. * - DEBUG_CLR_GRAPH_PACKET_CAPTURE - * 1) Added 2 nodes of MemCpy & a Kernel node in sequence and copy back result using memcpy + * 1) Added 2 nodes of MemCpy & a Kernel node in sequence and copy back result using memcpy and Instantiate graph & update new graph with similar node structure with api hipGraphExecUpdate and verify the result, the updated node should reflect. i) Check with Multi device case. diff --git a/catch/unit/graph/hipGraphRemoveDependencies.cc b/catch/unit/graph/hipGraphRemoveDependencies.cc index 2cbc28b749..44762cc9cc 100644 --- a/catch/unit/graph/hipGraphRemoveDependencies.cc +++ b/catch/unit/graph/hipGraphRemoveDependencies.cc @@ -19,8 +19,9 @@ THE SOFTWARE. #include #include + #include - + #include "graph_dependency_common.hh" @@ -307,6 +308,7 @@ TEST_CASE("Unit_hipGraphRemoveDependencies_Positive_ChangeComputeFunc") { REQUIRE(false == bMismatch); // Destroy HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } diff --git a/catch/unit/graph/hipGraphRemoveDependencies_old.cc b/catch/unit/graph/hipGraphRemoveDependencies_old.cc index 18cfd728cd..e41dc9e072 100644 --- a/catch/unit/graph/hipGraphRemoveDependencies_old.cc +++ b/catch/unit/graph/hipGraphRemoveDependencies_old.cc @@ -372,6 +372,7 @@ TEST_CASE("Unit_hipGraphRemoveDependencies_ChangeComputeFunc") { REQUIRE(false == bMismatch); // Destroy HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } @@ -459,4 +460,4 @@ TEST_CASE("Unit_hipGraphRemoveDependencies_Negative") { HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipEventDestroy(event_end)); HIP_CHECK(hipEventDestroy(event_start)); -} \ No newline at end of file +} diff --git a/catch/unit/graph/hipSimpleGraphWithKernel.cc b/catch/unit/graph/hipSimpleGraphWithKernel.cc index 08f20cbe45..537791d4a9 100644 --- a/catch/unit/graph/hipSimpleGraphWithKernel.cc +++ b/catch/unit/graph/hipSimpleGraphWithKernel.cc @@ -94,6 +94,8 @@ static void hipTestWithGraph() { delete[] out_h; HIP_CHECK(hipFree(in_d)); HIP_CHECK(hipFree(out_d)); + HIP_CHECK(hipGraphExecDestroy(instance)); + HIP_CHECK(hipStreamDestroy(stream)); } static void hipTestWithoutGraph() { @@ -143,6 +145,7 @@ static void hipTestWithoutGraph() { delete[] out_h; HIP_CHECK(hipFree(in_d)); HIP_CHECK(hipFree(out_d)); + HIP_CHECK(hipStreamDestroy(stream)); } /** diff --git a/catch/unit/graph/hipStreamBeginCapture.cc b/catch/unit/graph/hipStreamBeginCapture.cc index 78204efac9..ad15f2f451 100644 --- a/catch/unit/graph/hipStreamBeginCapture.cc +++ b/catch/unit/graph/hipStreamBeginCapture.cc @@ -19,7 +19,7 @@ THE SOFTWARE. #include #include - + #include "stream_capture_common.hh" // NOLINT #pragma clang diagnostic ignored "-Wunused-variable" @@ -218,7 +218,6 @@ static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStrea EventsGuard events_guard(1); hipEvent_t event = events_guard[0]; - HIP_CHECK(hipEventCreate(&event)); HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); HIP_CHECK(hipEventRecord(event, stream1)); HIP_CHECK(hipStreamWaitEvent(stream2, event, 0)); diff --git a/catch/unit/graph/hipStreamBeginCaptureToGraph.cc b/catch/unit/graph/hipStreamBeginCaptureToGraph.cc index 20b1d5b13c..4b56eae8a8 100644 --- a/catch/unit/graph/hipStreamBeginCaptureToGraph.cc +++ b/catch/unit/graph/hipStreamBeginCaptureToGraph.cc @@ -969,7 +969,6 @@ TEST_CASE("Unit_hipStreamBeginCaptureToGraph_StateTesting") { HIP_CHECK(hipStreamCreate(&stream2)); HIP_CHECK(hipEventCreate(&e)); hipStreamCaptureStatus captureStatus = hipStreamCaptureStatusNone; - HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipGraphCreate(&graph, 0)); HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus)); REQUIRE(captureStatus == hipStreamCaptureStatusNone); diff --git a/catch/unit/graph/hipStreamGetCaptureInfo_old.cc b/catch/unit/graph/hipStreamGetCaptureInfo_old.cc index c4b0b4b5b1..e6833716e4 100644 --- a/catch/unit/graph/hipStreamGetCaptureInfo_old.cc +++ b/catch/unit/graph/hipStreamGetCaptureInfo_old.cc @@ -149,6 +149,7 @@ void validateStreamCaptureInfo(hipStream_t mstream) { HIP_CHECK(hipStreamSynchronize(streamForLaunch)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForLaunch)); HIP_CHECK(hipStreamDestroy(stream1)); @@ -610,6 +611,7 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo") { REQUIRE(C_h[i] == D_h[i]); } + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipStreamDestroy(streamForGraph)); @@ -618,4 +620,4 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo") { free(A_h); free(C_h); free(D_h); -} \ No newline at end of file +} diff --git a/catch/unit/graph/hipStreamGetCaptureInfo_v2_old.cc b/catch/unit/graph/hipStreamGetCaptureInfo_v2_old.cc index 47a168b02c..3d833d4250 100644 --- a/catch/unit/graph/hipStreamGetCaptureInfo_v2_old.cc +++ b/catch/unit/graph/hipStreamGetCaptureInfo_v2_old.cc @@ -158,6 +158,7 @@ void validateStreamCaptureInfoV2(hipStream_t mstream) { HIP_CHECK(hipStreamSynchronize(streamForLaunch)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForLaunch)); HIP_CHECK(hipStreamDestroy(stream1)); @@ -277,4 +278,4 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_ParamValidation") { HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipFree(A_d)); -} \ No newline at end of file +} diff --git a/catch/unit/graph/hipStreamIsCapturing_old.cc b/catch/unit/graph/hipStreamIsCapturing_old.cc index ba4634a394..d963d43f49 100644 --- a/catch/unit/graph/hipStreamIsCapturing_old.cc +++ b/catch/unit/graph/hipStreamIsCapturing_old.cc @@ -425,6 +425,7 @@ TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") { D_h[i] = A_h[i] * A_h[i]; REQUIRE(C_h[i] == D_h[i]); } + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipStreamDestroy(streamForGraph)); @@ -433,4 +434,4 @@ TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") { free(A_h); free(C_h); free(D_h); -} \ No newline at end of file +} diff --git a/catch/unit/memory/hipMallocManaged.cc b/catch/unit/memory/hipMallocManaged.cc index 1b9766ffa9..04f579f0fa 100644 --- a/catch/unit/memory/hipMallocManaged.cc +++ b/catch/unit/memory/hipMallocManaged.cc @@ -138,6 +138,9 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") { } HIP_CHECK(hipFree(A)); HIP_CHECK(hipFree(B)); + HIP_CHECK(hipFree(C)); + HIP_CHECK(hipEventDestroy(event0)); + HIP_CHECK(hipEventDestroy(event1)); REQUIRE(maxError != 0.0f); } diff --git a/catch/unit/memory/hipMallocMngdMultiThread.cc b/catch/unit/memory/hipMallocMngdMultiThread.cc index a163e9a0ab..4482c5e352 100644 --- a/catch/unit/memory/hipMallocMngdMultiThread.cc +++ b/catch/unit/memory/hipMallocMngdMultiThread.cc @@ -119,6 +119,9 @@ static void LaunchKrnl2(int* Hmm, size_t NumElms, int InitVal, int HmmMem) { INFO("Data Mismatch observed at line: " << __LINE__); REQUIRE(false); } + + HIP_CHECK(hipFree(ptr)); + HIP_CHECK(hipStreamDestroy(strm)); } static void LaunchKrnl3(int* Dptr, size_t NumElms, int InitVal) { @@ -143,6 +146,9 @@ static void LaunchKrnl3(int* Dptr, size_t NumElms, int InitVal) { INFO("Data Mismatch observed at line: " << __LINE__); REQUIRE(false); } + + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); } @@ -180,6 +186,9 @@ static void LaunchKrnl5(int* Hmm1, size_t NumElms, int InitVal, int KerneltoLaun INFO("Data Mismatch observed at line: " << __LINE__); REQUIRE(false); } + + HIP_CHECK(hipFree(Hmm2)); + HIP_CHECK(hipStreamDestroy(strm)); } diff --git a/catch/unit/memory/hipMemRangeGetAttributes_old.cc b/catch/unit/memory/hipMemRangeGetAttributes_old.cc index 63952f1bfd..7ca0f87c17 100644 --- a/catch/unit/memory/hipMemRangeGetAttributes_old.cc +++ b/catch/unit/memory/hipMemRangeGetAttributes_old.cc @@ -156,6 +156,7 @@ TEST_CASE("Unit_hipMemRangeGetAttributes_TstFlgs") { } HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); delete[] AcsdBy; for (int i = 0; i < 4; ++i) { delete Outpt[i]; diff --git a/catch/unit/memory/hipMemcpy2DAsync_old.cc b/catch/unit/memory/hipMemcpy2DAsync_old.cc index 233d6077dd..69f8c73b6a 100644 --- a/catch/unit/memory/hipMemcpy2DAsync_old.cc +++ b/catch/unit/memory/hipMemcpy2DAsync_old.cc @@ -418,6 +418,7 @@ TEST_CASE("Unit_hipMemcpy2DAsync_SizeCheck") { // DeAllocating the memory HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); free(A_h); } diff --git a/catch/unit/memory/hipMemset.cc b/catch/unit/memory/hipMemset.cc index c4adb15e53..20bde29041 100644 --- a/catch/unit/memory/hipMemset.cc +++ b/catch/unit/memory/hipMemset.cc @@ -278,4 +278,7 @@ TEST_CASE("Unit_hipMemset_2AsyncOperations") { REQUIRE(v[0] == 0); REQUIRE(v[1024] == 1.75f); + + HIP_CHECK(hipFree(p2)); + HIP_CHECK(hipStreamDestroy(s)); } diff --git a/catch/unit/module/hipExtModuleLaunchKernel.cc b/catch/unit/module/hipExtModuleLaunchKernel.cc index c1d64b7388..2b3e9bc447 100644 --- a/catch/unit/module/hipExtModuleLaunchKernel.cc +++ b/catch/unit/module/hipExtModuleLaunchKernel.cc @@ -254,6 +254,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") { start_event, nullptr)); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipEventQuery(start_event)); + HIP_CHECK(hipEventDestroy(start_event)); } SECTION("Pass only stop event") { @@ -265,6 +266,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") { nullptr, stop_event)); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipEventQuery(stop_event)); + HIP_CHECK(hipEventDestroy(stop_event)); } } diff --git a/catch/unit/multiThread/hipMultiThreadStreams2.cc b/catch/unit/multiThread/hipMultiThreadStreams2.cc index d38754724d..c0968d0f37 100644 --- a/catch/unit/multiThread/hipMultiThreadStreams2.cc +++ b/catch/unit/multiThread/hipMultiThreadStreams2.cc @@ -51,7 +51,7 @@ void run1(size_t size, hipStream_t stream) { HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream)); HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream)); hipLaunchKernelGGL(HIP_KERNEL_NAME(Inc), dim3(N / 500), dim3(500), 0, stream, Cd); - HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipGetLastError()); HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream)); HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream)); HIPCHECK(hipDeviceSynchronize()); @@ -93,9 +93,9 @@ void run(size_t size, hipStream_t stream1, hipStream_t stream2) { HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1)); HIPCHECK(hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2)); hipLaunchKernelGGL(HIP_KERNEL_NAME(Inc), dim3(N / 500), dim3(500), 0, stream1, Cd); - HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipGetLastError()); hipLaunchKernelGGL(HIP_KERNEL_NAME(Inc), dim3(N / 500), dim3(500), 0, stream2, Cdd); - HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipGetLastError()); HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1)); HIPCHECK(hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2)); HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1)); @@ -140,4 +140,8 @@ TEST_CASE("Unit_hipMultiThreadStreams2") { t2.join(); t3.join(); } + + for (int i = 0; i < 3; i++) { + HIPCHECK(hipStreamDestroy(stream[i])); + } } diff --git a/catch/unit/rtc/hipStreamCaptureRtc.cc b/catch/unit/rtc/hipStreamCaptureRtc.cc index ab5296c4d5..2264b4fa57 100644 --- a/catch/unit/rtc/hipStreamCaptureRtc.cc +++ b/catch/unit/rtc/hipStreamCaptureRtc.cc @@ -123,6 +123,7 @@ TEST_CASE("Unit_hipStreamCaptureRtc") { HIPCHECK(hipGraphLaunch(graph_exec, stream)); HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipGraphExecDestroy(graph_exec)); HIPCHECK(hipStreamDestroy(stream)); // Check that the work was done diff --git a/catch/unit/stream/hipStreamCreateWithPriority.cc b/catch/unit/stream/hipStreamCreateWithPriority.cc index 7a246dd84e..37370bf500 100644 --- a/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/catch/unit/stream/hipStreamCreateWithPriority.cc @@ -488,12 +488,16 @@ bool validateStreamPrioritiesWithEvents() { OP(low, high) #undef OP - // free host & device memory + // free host & device memory & events #define OP(x) \ free(src_h_##x); \ free(dst_h_##x); \ HIP_CHECK(hipFree(src_d_##x)); \ - HIP_CHECK(hipFree(dst_d_##x)); + HIP_CHECK(hipFree(dst_d_##x)); \ + if (enable_priority_##x) { \ + HIP_CHECK(hipEventDestroy(event_start_##x)); \ + HIP_CHECK(hipEventDestroy(event_end_##x)); \ + } OP(low) OP(normal) OP(high) @@ -801,6 +805,33 @@ void TestForMultipleStreamWithPriority(void) { REQUIRE(memcmp(dst_h_high[i], src_h_high[i], size) == 0); } } + + for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; i++) { + if (enable_priority_low) { + HIP_CHECK(hipEventDestroy(event_start_low[i])); + HIP_CHECK(hipEventDestroy(event_end_low[i])); + } + + HIP_CHECK(hipStreamDestroy(stream_low[i])); + } + + for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; i++) { + if (enable_priority_normal) { + HIP_CHECK(hipEventDestroy(event_start_normal[i])); + HIP_CHECK(hipEventDestroy(event_end_normal[i])); + } + + HIP_CHECK(hipStreamDestroy(stream_normal[i])); + } + + for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; i++) { + if (enable_priority_high) { + HIP_CHECK(hipEventDestroy(event_start_high[i])); + HIP_CHECK(hipEventDestroy(event_end_high[i])); + } + + HIP_CHECK(hipStreamDestroy(stream_high[i])); + } } } // namespace hipStreamCreateWithPriorityTest diff --git a/catch/unit/stream/hipStreamGetDevice.cc b/catch/unit/stream/hipStreamGetDevice.cc index fbe8324270..930fc1a0ea 100644 --- a/catch/unit/stream/hipStreamGetDevice.cc +++ b/catch/unit/stream/hipStreamGetDevice.cc @@ -140,6 +140,7 @@ static bool validateStreamGetDevice() { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipStreamGetDevice(stream, &device_from_stream)); + HIP_CHECK(hipStreamDestroy(stream)); REQUIRE(device_from_stream == gpu); return true; diff --git a/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index f940bf1766..d03a7dbae1 100644 --- a/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -251,6 +251,9 @@ static void EventSync() { } else { IfTestPassed = true; } + + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(end)); } /* Launch a kernel in hipStreamPerThread, while it is in flight check for @@ -475,6 +478,7 @@ TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") { } HIP_CHECK(hipFree(Ad)); HIP_CHECK(hipFree(Ad1)); + HIP_CHECK(hipEventDestroy(e1)); HIP_CHECK(hipStreamDestroy(Strm)); delete[] Ah; delete Ah1; diff --git a/catch/unit/streamperthread/hipStreamPerThread_Event.cc b/catch/unit/streamperthread/hipStreamPerThread_Event.cc index 6fef9914eb..1e721f2757 100644 --- a/catch/unit/streamperthread/hipStreamPerThread_Event.cc +++ b/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(hipEventDestroy(event)); } __global__ void update_even_odd(unsigned int N, int* out) { @@ -61,4 +62,9 @@ TEST_CASE("Unit_hipStreamPerThread_EventSynchronize") { REQUIRE(false); } } -} \ No newline at end of file + + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(end)); +}