From 01519a8d5154d7404137960ee43f9048ad10cc49 Mon Sep 17 00:00:00 2001
From: "Yendluri, Manas" <Manas.Yendluri@amd.com>
Date: Tue, 2 Apr 2024 00:56:55 +0530
Subject: [PATCH] SWDEV-450972 - [catch2][dtest] Basic functional testcases to
 trigger capturehipMempcyHtoDAsync capturehipMemcpyDtoHAsync internal APIs

Change-Id: I2604dddf1d120622d565a1b74c93550e27ace031


[ROCm/hip-tests commit: 7c9f7737c0ea45052144066a84a97cf7425db6c7]
---
 .../unit/memory/hipMemcpyAsync_derivatives.cc | 94 ++++++++++++++++++-
 1 file changed, 93 insertions(+), 1 deletion(-)

diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc
index 251c630d52..7fed6121fe 100644
--- a/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc
+++ b/projects/hip-tests/catch/unit/memory/hipMemcpyAsync_derivatives.cc
@@ -172,4 +172,96 @@ TEST_CASE("Unit_hipMemcpyDtoDAsync_Negative_Parameters") {
                                        InvalidStream()),
                     hipErrorContextIsDestroyed);
   }
-}
\ No newline at end of file
+}
+
+/**
+* Test Description
+* ------------------------
+*  - Basic functional testcase to trigger capturehipMemcpyDtoHAsync internal api
+*  to improve code coverage.
+* Test source
+* ------------------------
+*  - unit/memory/hipMemcpyAsync_derivatives.cc
+* Test requirements
+* ------------------------
+*  - HIP_VERSION >= 6.0
+*/
+TEST_CASE("Unit_hipMemcpyDtoHAsync_capturehipMemcpyDtoHAsync") {
+  hipGraph_t graph{nullptr};
+  hipGraphExec_t graphExec{nullptr};
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  int *A_h = reinterpret_cast<int*>(malloc(sizeof(int) * kPageSize));
+  int *B_h = reinterpret_cast<int*>(malloc(sizeof(int) * kPageSize));
+  int *A_d;
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), sizeof(int) * kPageSize));
+  for (int i = 0; i < kPageSize; i++) {
+    B_h[i] = i;
+  }
+  HIP_CHECK(hipMemcpyHtoD((hipDeviceptr_t)A_d, B_h, sizeof(int) * kPageSize));
+  // Start Capturing
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipMemcpyDtoHAsync(A_h, (hipDeviceptr_t)A_d, sizeof(int) * kPageSize, 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));
+  for (int i = 0; i < kPageSize; i++) {
+    REQUIRE(A_h[i] ==  B_h[i]);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec))
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+  free(B_h);
+}
+
+/**
+* Test Description
+* ------------------------
+*  - Basic functional testcase to trigger capturehipMemcpyHtoDAsync internal api
+*  to improve code coverage.
+* Test source
+* ------------------------
+*  - unit/memory/hipMemcpyAsync_derivatives.cc
+* Test requirements
+* ------------------------
+*  - HIP_VERSION >= 6.0
+*/
+TEST_CASE("Unit_hipMemcpyHtoDAsync_capturehipMemcpyHtoDAsync") {
+  hipGraph_t graph{nullptr};
+  hipGraphExec_t graphExec{nullptr};
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  int *A_h = reinterpret_cast<int*>(malloc(sizeof(int) * kPageSize));
+  int *B_h = reinterpret_cast<int*>(malloc(sizeof(int) * kPageSize));
+  int *A_d;
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), sizeof(int) * kPageSize));
+  for (int i = 0; i < kPageSize; i++) {
+    B_h[i] = i;
+  }
+  // Start Capturing
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipMemcpyHtoDAsync((hipDeviceptr_t)A_d, B_h, sizeof(int) * kPageSize, 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(hipMemcpyDtoH(A_h, (hipDeviceptr_t)A_d, sizeof(int) * kPageSize));
+  for (int i = 0; i < kPageSize; i++) {
+    REQUIRE(A_h[i] == B_h[i]);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec))
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+  free(B_h);
+}