SWDEV-548417 - Fix Memleaks in Graph (#713)

Co-authored-by: Anusha GodavarthySurya <Anusha.GodavarthySurya@amd.com>

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);

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);
 

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, &copyParams_, 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);

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);

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();
 

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));

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<int>(A_d, B_d, C_d, A_h, B_h, C_h, false);
 }

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));

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));

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));
 }

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));

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<int>(A_d, B_d, C_d, A_h, B_h, nullptr, false);

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);
 }
 

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));
 }
 
 /**

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));
 }

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));
 }
 
 /**

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);

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<int*>(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

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));
 }

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));
 }

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));
 }