From d016f96947fbf6e850a3e7af8bd6b94bc48ec721 Mon Sep 17 00:00:00 2001
From: "Yendluri, Manas" <Manas.Yendluri@amd.com>
Date: Tue, 2 Apr 2024 01:01:44 +0530
Subject: [PATCH] SWDEV-450972 - [catch2][dtest] Basic functional testcase to
 trigger capturehipExtLaunchKernel internal api

Change-Id: I707bfe65de552f1831a66f9db3d60cb803ae6b3d


[ROCm/hip-tests commit: 8bd7575356ce048ac1e66534bdd0d9f7dfcc44da]
---
 .../executionControl/hipExtLaunchKernel.cc    | 48 ++++++++++++++++++-
 1 file changed, 47 insertions(+), 1 deletion(-)

diff --git a/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc
index 1b336b4d74..c0b4206733 100644
--- a/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc
+++ b/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc
@@ -173,4 +173,50 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") {
                                        dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, event, 0u),
                     hipErrorInvalidValue);
   }
-}
\ No newline at end of file
+}
+
+/**
+* Test Description
+* ------------------------
+*  - Basic functional testcase to trigger capturehipExtLaunchKernel internal api
+*  to improve code coverage.
+* Test source
+* ------------------------
+*  - unit/executionControl/hipExtLaunchKernel.cc
+* Test requirements
+* ------------------------
+*  - HIP_VERSION >= 6.0
+*/
+TEST_CASE("Unit_hipExtLaunchKernel_capturehipExtLaunchKernel") {
+  hipGraph_t graph{nullptr};
+  hipGraphExec_t graphExec{nullptr};
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  int *A_d;
+  int *A_h = nullptr;
+  A_h = reinterpret_cast<int*>(malloc(sizeof(int)));
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), sizeof(int)));
+  void *args[1] = {&A_d};
+  // Begin Capture operation
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipExtLaunchKernel(reinterpret_cast<void*>(kernel_42),
+                               dim3{1, 1, 1}, dim3{1, 1, 1}, args, 0, stream,
+                               nullptr, nullptr, 0u));
+  // 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(hipMemcpyDtoH(A_h, A_d, sizeof(int)));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(*A_h == 42);
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+}