SWDEV-450972 - [catch2][dtest] Basic functional testcases to trigger capturehipMempcyFromSymbolAsync capturehipMemcpyToSymbolAsync internal APIs

Change-Id: I8ade6f8f40e7a105792f07b1913a6b3fff477bd7


[ROCm/hip-tests commit: 6125f90e67]

projects/hip-tests/catch/unit/memory/hipMemcpyFromSymbol.cc

@@ -20,6 +20,7 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 
 __device__ int devSymbol[10];
+__constant__ int constSymbol[10];
 
 /* Test verifies hipMemcpy[From/To]Symbol[Async] API Negative scenarios.
  */
@@ -188,3 +189,51 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
     }
   }
 }
+
+/**
+* Test Description
+* ------------------------
+*  - Basic functional testcase to trigger capturehipMemcpyToSymbolAsync
+*  and capturehipMemcpyFromSymbolAsync internal apis to improve
+*  code coverage.
+* Test source
+* ------------------------
+*  - unit/memory/hipMemcpyFromSymbol.cc
+* Test requirements
+* ------------------------
+*  - HIP_VERSION >= 6.0
+*/
+TEST_CASE("Unit_hipMemcpyToFromSymbol_capturehipMemcpyToFromSymbolAsync") {
+  hipGraph_t graph{nullptr};
+  hipGraphExec_t graphExec{nullptr};
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+  int A_h = 0, B_h = 42;
+
+  // Start Capturing
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
+  SECTION("__constant__ symbol") {
+    HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(constSymbol), &B_h,
+              sizeof(int), 0, hipMemcpyHostToDevice, stream));
+    HIP_CHECK(hipMemcpyFromSymbolAsync(&A_h, HIP_SYMBOL(constSymbol),
+              sizeof(int), 0, hipMemcpyDeviceToHost, stream));
+  }
+  SECTION("__device__ symbol") {
+    HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &B_h,
+              sizeof(int), 0, hipMemcpyHostToDevice, stream));
+    HIP_CHECK(hipMemcpyFromSymbolAsync(&A_h, HIP_SYMBOL(devSymbol),
+              sizeof(int), 0, hipMemcpyDeviceToHost, 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));
+
+  REQUIRE(A_h == B_h);
+  HIP_CHECK(hipGraphExecDestroy(graphExec))
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream));
+}