From daba7f6f80c97ee865f0e1409272a2dd1a386775 Mon Sep 17 00:00:00 2001
From: Anusha GodavarthySurya <myendlur@amd.com>
Date: Tue, 2 Apr 2024 00:40:54 +0530
Subject: [PATCH] SWDEV-450972 - [catch2][dtest] Basic functional testcase to
 trigger capturehipFreeAsync internal API

Change-Id: I0e6fab330b9d91558842c6c81215bcec42c07220
---
 catch/unit/memory/hipFreeAsync.cc | 43 +++++++++++++++++++++++++++++++
 1 file changed, 43 insertions(+)

diff --git a/catch/unit/memory/hipFreeAsync.cc b/catch/unit/memory/hipFreeAsync.cc
index 7c49e1b76c..773433aead 100644
--- a/catch/unit/memory/hipFreeAsync.cc
+++ b/catch/unit/memory/hipFreeAsync.cc
@@ -82,3 +82,46 @@ TEST_CASE("Unit_hipFreeAsync_Negative_Parameters") {
 * End doxygen group StreamOTest.
 * @}
 */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Functional test cases to trigger capturehipFreeAsync internal api
+ * Test source
+ * ------------------------
+ *  - unit/memory/hipFreeAsync
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.0
+ */
+TEST_CASE("Unit_hipFreeAsync_capturehipFreeAsync") {
+  HIP_CHECK(hipSetDevice(0));
+  hipGraph_t graph{nullptr};
+  hipGraphExec_t graphExec{nullptr};
+  hipStream_t stream;
+  hipMemPool_t memPool;
+  int rows, cols;
+  rows = GENERATE(3, 4, 1024);
+  cols = GENERATE(3, 4, 1024);
+  HIP_CHECK(hipDeviceGetDefaultMemPool(&memPool, 0));
+  HIP_CHECK(hipStreamCreate(&stream));
+  int* devMem;
+
+  // Start Capturing
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&devMem),
+                                   sizeof(int) * rows * cols, memPool,
+                                   stream));
+  HIP_CHECK(hipFreeAsync(devMem, stream));
+  // End Capture
+  HIP_CHECK(hipStreamEndCapture(stream, &graph));
+
+  // Create and Launch Executable Graphs
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, stream));
+  HIP_CHECK(hipStreamSynchronize(stream));
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream));
+}