From ce560304a8e4600b001736cc8a685eede9e18c56 Mon Sep 17 00:00:00 2001 From: "Godavarthy Surya, Anusha" Date: Fri, 19 Sep 2025 17:39:36 +0530 Subject: [PATCH] SWDEV-548417 - Fix Memleaks in Graph (#713) Co-authored-by: Anusha GodavarthySurya --- projects/clr/hipamd/src/hip_graph.cpp | 1 + projects/clr/hipamd/src/hip_graph_helper.hpp | 2 +- .../clr/hipamd/src/hip_graph_internal.hpp | 18 ++++++---- projects/clr/hipamd/src/hip_memory.cpp | 2 +- projects/clr/rocclr/platform/commandqueue.cpp | 22 ++++++------- .../graph/hipDrvGraphMemcpyNodeGetParams.cc | 1 + .../catch/unit/graph/hipGraphAddKernelNode.cc | 1 + .../unit/graph/hipGraphAddNodeBeginCapture.cc | 4 +-- .../catch/unit/graph/hipGraphCloneComplx.cc | 6 ++++ .../hipGraphExecChildGraphNodeSetParams.cc | 1 + .../catch/unit/graph/hipGraphExecDestroy.cc | 1 + ...raphExecMemcpyNodeSetParamsToSymbol_old.cc | 1 + .../unit/graph/hipGraphNodeGetEnabled.cc | 1 + .../catch/unit/graph/hipGraphPerf.cc | 6 ++++ .../unit/graph/hipGraphRetainUserObject.cc | 33 ++----------------- .../catch/unit/graph/hipLaunchHostFunc.cc | 4 ++- .../catch/unit/graph/hipStreamEndCapture.cc | 2 +- .../unit/graph/hipStreamEndCapture_old.cc | 11 +++++-- .../catch/unit/graph/hipUserObjectCreate.cc | 14 ++++++-- .../catch/unit/graph/hipUserObjectRelease.cc | 1 + .../catch/unit/graph/hipUserObjectRetain.cc | 6 +++- 21 files changed, 78 insertions(+), 60 deletions(-) diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 878f2dfd11..b730605e29 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -1518,6 +1518,7 @@ hipError_t ihipGraphInstantiate(hip::GraphExec** pGraphExec, hip::Graph* graph, graph->clone(*pGraphExec, true); (*pGraphExec)->ScheduleNodes(); if (false == (*pGraphExec)->TopologicalOrder()) { + delete *pGraphExec; return hipErrorInvalidValue; } graph->SetGraphInstantiated(true); diff --git a/projects/clr/hipamd/src/hip_graph_helper.hpp b/projects/clr/hipamd/src/hip_graph_helper.hpp index 0b79d3fcb4..f157f4c8c9 100644 --- a/projects/clr/hipamd/src/hip_graph_helper.hpp +++ b/projects/clr/hipamd/src/hip_graph_helper.hpp @@ -35,7 +35,7 @@ hipError_t hipMemcpy2DValidateBuffer(const void* buf, size_t pitch, size_t width hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind); hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes, - hipMemcpyKind kind, hip::Stream& stream, bool isAsync = false); + hipMemcpyKind kind, hip::Stream& stream, bool isAsync = true); void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, hip::Stream& stream); diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index c6767508c9..107bcd95d4 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -1241,6 +1241,9 @@ class GraphKernelNode : public GraphNode { GraphNode* clone() const override { return new GraphKernelNode(*this); } hipError_t CreateCommand(hip::Stream* stream) override { + if (!isEnabled_) { + return hipSuccess; + } hipFunction_t func = getFunc(kernelParams_, dev_id_); if (!func) { return hipErrorInvalidDeviceFunction; @@ -1468,8 +1471,9 @@ class GraphMemcpyNode : public GraphNode { GraphNode* clone() const override { return new GraphMemcpyNode(*this); } virtual hipError_t CreateCommand(hip::Stream* stream) override { - if ((copyParams_.kind == hipMemcpyHostToHost || copyParams_.kind == hipMemcpyDefault) && - IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr)) { + if (!isEnabled_ || + ((copyParams_.kind == hipMemcpyHostToHost || copyParams_.kind == hipMemcpyDefault) && + IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr))) { return hipSuccess; } hipError_t status = GraphNode::CreateCommand(stream); @@ -1480,6 +1484,7 @@ class GraphMemcpyNode : public GraphNode { amd::Command* command; status = ihipMemcpy3DCommand(command, ©Params_, stream); commands_.emplace_back(command); + return status; } @@ -1671,7 +1676,8 @@ class GraphMemcpyNode1D : public GraphMemcpyNode { GraphNode* clone() const override { return new GraphMemcpyNode1D(*this); } virtual hipError_t CreateCommand(hip::Stream* stream) override { - if ((kind_ == hipMemcpyHostToHost || kind_ == hipMemcpyDefault) && IsHtoHMemcpy(dst_, src_)) { + if (!isEnabled_ || + ((kind_ == hipMemcpyHostToHost || kind_ == hipMemcpyDefault) && IsHtoHMemcpy(dst_, src_))) { return hipSuccess; } hipError_t status = GraphNode::CreateCommand(stream); @@ -2610,9 +2616,9 @@ class GraphDrvMemcpyNode : public GraphNode { GraphNode* clone() const override { return new GraphDrvMemcpyNode(*this); } hipError_t CreateCommand(hip::Stream* stream) override { - if (copyParams_.srcMemoryType == hipMemoryTypeHost && - copyParams_.dstMemoryType == hipMemoryTypeHost && - IsHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost)) { + if (!isEnabled_ || (copyParams_.srcMemoryType == hipMemoryTypeHost && + copyParams_.dstMemoryType == hipMemoryTypeHost && + IsHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost))) { return hipSuccess; } hipError_t status = GraphNode::CreateCommand(stream); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 5cebd56a8f..0cfd93f77a 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1795,7 +1795,7 @@ hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* dstHost, amd::Coo } else { amd::Command::EventWaitList waitList; auto* pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); - if (stream != pStream) { + if (stream->DeviceId() != srcMemory->getUserData().deviceId) { amd::Command* cmd = pStream->getLastQueuedCommand(true); if (cmd != nullptr) { waitList.push_back(cmd); diff --git a/projects/clr/rocclr/platform/commandqueue.cpp b/projects/clr/rocclr/platform/commandqueue.cpp index b965c9fad1..c6f96b6530 100644 --- a/projects/clr/rocclr/platform/commandqueue.cpp +++ b/projects/clr/rocclr/platform/commandqueue.cpp @@ -208,20 +208,18 @@ void HostQueue::finish(bool cpu_wait) { "await command completion", minBatchSize); command->awaitCompletion(); - - if (IS_HIP) { - ScopedLock sl(vdev()->execution()); - ScopedLock l(lastCmdLock_); - // Runtime can clear the last command only if no other submissions occured - // during finish() - if (command == lastEnqueueCommand_) { - device_.removeFromActiveQueues(this); - lastEnqueueCommand_->release(); - lastEnqueueCommand_ = nullptr; - } + } + if (IS_HIP) { + ScopedLock sl(vdev()->execution()); + ScopedLock l(lastCmdLock_); + // Runtime can clear the last command only if no other submissions occured + // during finish() + if (command == lastEnqueueCommand_) { + device_.removeFromActiveQueues(this); + lastEnqueueCommand_->release(); + lastEnqueueCommand_ = nullptr; } } - // Release all HW queues, which are idle or nearly idle vdev()->ReleaseAllHwQueues(); diff --git a/projects/hip-tests/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc b/projects/hip-tests/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc index c065278ae7..77c9985a6f 100644 --- a/projects/hip-tests/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc @@ -174,6 +174,7 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeGetParams_Positive") { REQUIRE(memCpy_params.dstHeight == memCpyGetParams.dstHeight); HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipStreamDestroy(streamForGraph)); HIP_CHECK(hipCtxPopCurrent(&context)); HIP_CHECK(hipCtxDestroy(context)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddKernelNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddKernelNode.cc index ef09055afe..a8e99688a7 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddKernelNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddKernelNode.cc @@ -393,6 +393,7 @@ TEST_CASE("Unit_hipGraphAddKernelNode_moduleLoadKernelFn_childGraph") { HIP_CHECK(hipGraphDestroy(clonedGraph)); } HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(childgraph)); HIPCHECK(hipModuleUnload(Module)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddNodeBeginCapture.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddNodeBeginCapture.cc index 6bc8d4f642..c1c7fb5ac3 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddNodeBeginCapture.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddNodeBeginCapture.cc @@ -99,7 +99,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_with_hipGraphAddHostNode") { p.fn = CpuCallback; p.userData = nullptr; HIP_CHECK(hipGraphAddHostNode(&cpuGraphNode, graph, nullptr, 0, &p)); - + HIP_CHECK(hipGraphDestroy(graph)); addGpuKernel<<<1, 1, 0, stream>>>(i_d); HIP_CHECK(hipStreamEndCapture(stream, &graph)); @@ -795,8 +795,6 @@ TEST_CASE("Unit_hipStreamEndCapture_first_and_add_other_graph_node_later") { HipTest::initArrays(&A_d1, &B_d1, &C_d1, &A_h1, &B_h1, &C_h1, SIZE, false); HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); HIP_CHECK(hipMemcpyAsync(A_d1, A_h1, Nbytes, hipMemcpyHostToDevice, stream)); HIP_CHECK(hipMemcpyAsync(B_d1, B_h1, Nbytes, hipMemcpyHostToDevice, stream)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphCloneComplx.cc b/projects/hip-tests/catch/unit/graph/hipGraphCloneComplx.cc index d8857ab704..ada6801727 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphCloneComplx.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphCloneComplx.cc @@ -642,6 +642,7 @@ static void hipGraphClone_Test_hipGraphMemcpyNodeSetParams1D_and_exec() { HIP_CHECK(hipGraphExecDestroy(clonedGraphExec_2)); } HIP_CHECK(hipGraphDestroy(clonedGraph)); + HIP_CHECK(hipGraphDestroy(clonedGraph_2)); } TEST_CASE("Unit_hipGraphClone_Test_hipGraphMemcpyNodeSetParams1D_and_exec") { @@ -735,6 +736,7 @@ static void hipGraphClone_hipGraphMemcpyNodeSetParamsFromSymbol_exec() { HIP_CHECK(hipGraphExecDestroy(clonedGraphExec_2)); } HIP_CHECK(hipGraphDestroy(clonedGraph)); + HIP_CHECK(hipGraphDestroy(clonedGraph_2)); } TEST_CASE("Unit_hipGraphClone_hipGraphMemcpyNodeSetParamsFromSymbol_exec") { @@ -827,6 +829,7 @@ static void hipGraphClone_hipGraphMemcpyNodeSetParamsToSymbol_exec() { HIP_CHECK(hipGraphExecDestroy(clonedGraphExec_2)); } HIP_CHECK(hipGraphDestroy(clonedGraph)); + HIP_CHECK(hipGraphDestroy(clonedGraph_2)); } TEST_CASE("Unit_hipGraphClone_hipGraphMemcpyNodeSetParamsToSymbol_exec") { @@ -938,6 +941,7 @@ static void hipGraphClone_Test_hipGraphMemsetNodeSetParams_exec() { HIP_CHECK(hipGraphExecDestroy(clonedGraphExec_2)); } HIP_CHECK(hipGraphDestroy(clonedGraph)); + HIP_CHECK(hipGraphDestroy(clonedGraph_2)); } TEST_CASE("Unit_hipGraphClone_Test_hipGraphMemsetNodeSetParams_exec") { @@ -1295,6 +1299,8 @@ static void hipGraphClone_Test_hipGraphEventRecordNodeSetEvent_and_Exec() { } HIP_CHECK(hipGraphDestroy(clonedGraph)); + HIP_CHECK(hipGraphDestroy(clonedGraph_3)); + HIP_CHECK(hipGraphDestroy(clonedGraph_4)); HIP_CHECK(hipGraphDestroy(childgraph)); HIP_CHECK(hipEventDestroy(event_start)); HIP_CHECK(hipEventDestroy(event_end)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc index 3f00ddec53..37d0ca6755 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc @@ -305,6 +305,7 @@ TEST_CASE("Unit_hipGraphExecChildGraphNodeSetParams_ChildTopology") { HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(childgraph1)); + HIP_CHECK(hipGraphDestroy(childgraph2)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc index f45f72cea6..f0c8b913d9 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc @@ -89,6 +89,7 @@ TEST_CASE("Unit_hipGraphExecDestroy_Positive_Basic") { HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); HIP_CHECK(hipStreamSynchronize(streamForGraph)); + HIP_CHECK(hipFree(devData)); HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipStreamDestroy(streamForGraph)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol_old.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol_old.cc index bcb2e920aa..1c14f524f3 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol_old.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol_old.cc @@ -151,6 +151,7 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative") { HIP_SYMBOL(globalOut), HIP_SYMBOL(globalIn), Nbytes, 0, hipMemcpyDeviceToDevice); REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK(hipGraphDestroy(graph1)); } HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, nullptr, false); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphNodeGetEnabled.cc b/projects/hip-tests/catch/unit/graph/hipGraphNodeGetEnabled.cc index 9faa6d9de7..e8c96684de 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphNodeGetEnabled.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphNodeGetEnabled.cc @@ -139,6 +139,7 @@ TEST_CASE("Unit_hipGraphNodeGetEnabled_Functional_Basic") { HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(childGraph)); HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphPerf.cc b/projects/hip-tests/catch/unit/graph/hipGraphPerf.cc index 5241ac664b..990c0c8216 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphPerf.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphPerf.cc @@ -1569,6 +1569,8 @@ static void hipGraph_PerfCheck_hipGraphExecChildGraphNodeSetParams(const hipStre 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(hipGraphDestroy(childGraph1)); + HIP_CHECK(hipGraphDestroy(childGraph2)); } static void hipGraph_PerfCheck_hipGraphExecChildGraphNodeSetParams_Kernel( @@ -1688,6 +1690,8 @@ static void hipGraph_PerfCheck_hipGraphExecChildGraphNodeSetParams_Kernel( HipTest::freeArrays(A_d1, B_d1, C_d1, A_h1, B_h1, C_h1, false); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(childGraph1)); + HIP_CHECK(hipGraphDestroy(childGraph2)); } static void hipGraph_PerfCheck_hipGraphExecChildGraphNodeSetParams_mKernel( @@ -1824,6 +1828,8 @@ static void hipGraph_PerfCheck_hipGraphExecChildGraphNodeSetParams_mKernel( HipTest::freeArrays(A_d1, B_d1, C_d1, A_h1, B_h1, C_h1, false); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(childGraph1)); + HIP_CHECK(hipGraphDestroy(childGraph2)); } /** diff --git a/projects/hip-tests/catch/unit/graph/hipGraphRetainUserObject.cc b/projects/hip-tests/catch/unit/graph/hipGraphRetainUserObject.cc index ebbe322d17..f04b026b01 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphRetainUserObject.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphRetainUserObject.cc @@ -208,36 +208,9 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Basic") { // Again Retain graph object with reference count 8 HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 8, hipGraphUserObjectMove)); - // Release graph object with reference count 1 - HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 1)); + // Release graph object with reference count 8 + HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 8)); - HIP_CHECK(hipUserObjectRelease(hObject, 1)); - HIP_CHECK(hipGraphDestroy(graph)); -} - -TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Null_Object") { - hipGraph_t graph; - HIP_CHECK(hipGraphCreate(&graph, 0)); - - float* object = nullptr; // this is used for Null_Object test - hipUserObject_t hObject; - - HIP_CHECK( - hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync)); - REQUIRE(hObject != nullptr); - - // Retain graph object with reference count 2 - HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2, hipGraphUserObjectMove)); - - // Release graph object with reference count more than 2 - HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 4)); - - // Again Retain graph object with reference count 8 - HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 8, 0)); - - // Release graph object with reference count 1 - HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, 1)); - - HIP_CHECK(hipUserObjectRelease(hObject, 1)); + HIP_CHECK(hipUserObjectRelease(hObject, 2)); HIP_CHECK(hipGraphDestroy(graph)); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc b/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc index 90ac633c2b..49ee0c867c 100644 --- a/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc +++ b/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc @@ -58,7 +58,7 @@ static void hostNodeCallback(void* data) { TEST_CASE("Unit_hipLaunchHostFunc_Negative_Parameters") { StreamGuard stream_guard(Streams::created); hipStream_t stream = stream_guard.stream(); - + hipGraph_t graph{nullptr}; HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); #if HT_NVIDIA // EXSWHTEC-228 SECTION("Pass stream as nullptr") { @@ -69,6 +69,8 @@ TEST_CASE("Unit_hipLaunchHostFunc_Negative_Parameters") { SECTION("Pass functions as nullptr") { HIP_CHECK_ERROR(hipLaunchHostFunc(stream, nullptr, nullptr), hipErrorInvalidValue); } + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipGraphDestroy(graph)); } /** diff --git a/projects/hip-tests/catch/unit/graph/hipStreamEndCapture.cc b/projects/hip-tests/catch/unit/graph/hipStreamEndCapture.cc index 18df4c55b4..49c8bd116e 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamEndCapture.cc +++ b/projects/hip-tests/catch/unit/graph/hipStreamEndCapture.cc @@ -96,6 +96,7 @@ TEST_CASE("Unit_hipStreamEndCapture_Positive_GraphDestroy") { HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipGraphDestroy(graph)); } static void thread_func_neg(hipStream_t stream, hipGraph_t graph) { @@ -128,7 +129,6 @@ TEST_CASE("Unit_hipStreamEndCapture_Negative_Thread") { hipStream_t stream = stream_guard.stream(); const hipStreamCaptureMode captureMode = hipStreamCaptureModeGlobal; - HIP_CHECK(hipGraphCreate(&graph, 0)); HIP_CHECK(hipStreamBeginCapture(stream, captureMode)); captureSequenceSimple(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), N, stream); diff --git a/projects/hip-tests/catch/unit/graph/hipStreamEndCapture_old.cc b/projects/hip-tests/catch/unit/graph/hipStreamEndCapture_old.cc index 529737151d..7d3c1b6828 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamEndCapture_old.cc +++ b/projects/hip-tests/catch/unit/graph/hipStreamEndCapture_old.cc @@ -124,6 +124,7 @@ TEST_CASE("Unit_hipStreamEndCapture_Negative") { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(C_d)); HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipGraphDestroy(graph)); } } @@ -140,7 +141,6 @@ static void StreamEndCaptureThreadNegative(float* A_d, float* A_h, float* C_d, f size_t Nbytes = N * sizeof(float); HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipGraphCreate(&graph, 0)); HIP_CHECK(hipStreamBeginCapture(stream, mode)); HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); @@ -277,6 +277,7 @@ TEST_CASE("Unit_hipStreamEndCapture_chkError_on_wrongStream") { int *A_d{nullptr}, *A_h{nullptr}; hipStream_t stream1{nullptr}, stream2{nullptr}; hipEvent_t forkStreamEvent{nullptr}; + hipEvent_t joinStreamEvent{nullptr}; hipGraph_t graph{nullptr}; hipError_t err; constexpr unsigned blocks = 512; @@ -286,7 +287,7 @@ TEST_CASE("Unit_hipStreamEndCapture_chkError_on_wrongStream") { HIP_CHECK(hipStreamCreate(&stream1)); HIP_CHECK(hipStreamCreate(&stream2)); HIP_CHECK(hipEventCreate(&forkStreamEvent)); - + HIP_CHECK(hipEventCreate(&joinStreamEvent)); A_h = reinterpret_cast(malloc(Nbytes)); REQUIRE(A_h != nullptr); // Initialize the Host data @@ -305,10 +306,15 @@ TEST_CASE("Unit_hipStreamEndCapture_chkError_on_wrongStream") { err = hipStreamEndCapture(stream2, &graph); REQUIRE(err == hipErrorStreamCaptureUnmatched); + HIP_CHECK(hipEventRecord(joinStreamEvent, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream1, joinStreamEvent, 0)); + err = hipStreamEndCapture(stream1, &graph); HIP_CHECK(hipStreamDestroy(stream1)); HIP_CHECK(hipStreamDestroy(stream2)); HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(joinStreamEvent)); + HIP_CHECK(hipGraphDestroy(graph)); free(A_h); HIP_CHECK(hipFree(A_d)); } @@ -404,6 +410,7 @@ TEST_CASE("Unit_hipStreamEndCapture_streamMerge_in_thread") { HIP_CHECK(hipStreamDestroy(stream1)); HIP_CHECK(hipStreamDestroy(stream2)); HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(event)); HIP_CHECK(hipStreamDestroy(streamForGraph)); // Release the memory diff --git a/projects/hip-tests/catch/unit/graph/hipUserObjectCreate.cc b/projects/hip-tests/catch/unit/graph/hipUserObjectCreate.cc index 2a22e62d9f..f8e781a1ea 100644 --- a/projects/hip-tests/catch/unit/graph/hipUserObjectCreate.cc +++ b/projects/hip-tests/catch/unit/graph/hipUserObjectCreate.cc @@ -190,28 +190,35 @@ TEST_CASE("Unit_hipUserObjectCreate_Negative") { HIP_CHECK_ERROR( hipUserObjectCreate(nullptr, object, destroyIntObj, 1, hipUserObjectNoDestructorSync), hipErrorInvalidValue); + delete object; } SECTION("Pass object as nullptr") { HIP_CHECK( hipUserObjectCreate(&hObject, nullptr, destroyIntObj, 1, hipUserObjectNoDestructorSync)); + HIP_CHECK(hipUserObjectRelease(hObject, 1)); + delete object; } SECTION("Pass Callback function as nullptr") { HIP_CHECK_ERROR( hipUserObjectCreate(&hObject, object, nullptr, 1, hipUserObjectNoDestructorSync), hipErrorInvalidValue); + delete object; } SECTION("Pass initialRefcount as 0") { HIP_CHECK_ERROR( hipUserObjectCreate(&hObject, object, destroyIntObj, 0, hipUserObjectNoDestructorSync), hipErrorInvalidValue); + delete object; } SECTION("Pass initialRefcount as INT_MAX") { HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyIntObj, INT_MAX, hipUserObjectNoDestructorSync)); + HIP_CHECK(hipUserObjectRelease(hObject, INT_MAX)); } SECTION("Pass flag other than hipUserObjectNoDestructorSync") { HIP_CHECK_ERROR(hipUserObjectCreate(&hObject, object, destroyIntObj, 1, hipUserObjectFlags(9)), hipErrorInvalidValue); + delete object; } } @@ -225,9 +232,12 @@ TEST_CASE("Unit_hipUserObj_Negative_Test") { HIP_CHECK(hipUserObjectCreate(&hObject, object, destroyIntObj, 2, hipUserObjectNoDestructorSync)); REQUIRE(hObject != nullptr); - // Release more than created. + // Release more than created. Will not release the references. HIP_CHECK(hipUserObjectRelease(hObject, 4)); - // Retain reference to a removed user object + // Retain reference HIP_CHECK(hipUserObjectRetain(hObject, 1)); + + // Release all the references pending + HIP_CHECK(hipUserObjectRelease(hObject, 3)); } diff --git a/projects/hip-tests/catch/unit/graph/hipUserObjectRelease.cc b/projects/hip-tests/catch/unit/graph/hipUserObjectRelease.cc index bb73fdeac2..0b56062967 100644 --- a/projects/hip-tests/catch/unit/graph/hipUserObjectRelease.cc +++ b/projects/hip-tests/catch/unit/graph/hipUserObjectRelease.cc @@ -45,4 +45,5 @@ TEST_CASE("Unit_hipUserObjectRelease_Negative") { HIP_CHECK_ERROR(hipUserObjectRelease(hObject, 0), hipErrorInvalidValue); } SECTION("Pass initialRefcount as INT_MAX") { HIP_CHECK(hipUserObjectRelease(hObject, INT_MAX)); } + HIP_CHECK(hipUserObjectRelease(hObject, 1)); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/graph/hipUserObjectRetain.cc b/projects/hip-tests/catch/unit/graph/hipUserObjectRetain.cc index 895c898482..5fb48f10c6 100644 --- a/projects/hip-tests/catch/unit/graph/hipUserObjectRetain.cc +++ b/projects/hip-tests/catch/unit/graph/hipUserObjectRetain.cc @@ -44,5 +44,9 @@ TEST_CASE("Unit_hipUserObjectRetain_Negative") { SECTION("Pass initialRefcount as 0") { HIP_CHECK_ERROR(hipUserObjectRetain(hObject, 0), hipErrorInvalidValue); } - SECTION("Pass initialRefcount as INT_MAX") { HIP_CHECK(hipUserObjectRetain(hObject, INT_MAX)); } + SECTION("Pass initialRefcount as INT_MAX") { + HIP_CHECK(hipUserObjectRetain(hObject, INT_MAX)); + HIP_CHECK(hipUserObjectRelease(hObject, INT_MAX)); + } + HIP_CHECK(hipUserObjectRelease(hObject, 1)); } \ No newline at end of file