SWDEV-306122 - [catch2][dtest] hipGraph tests for hipGraphAddMemcpyNode api (#2498)

Change-Id: I65088f4107158e320d54b5e54a47d4ad8ce5d85f
2022-02-24 18:57:27 +05:30
parent d305e34173
commit 7d8d41e45b
1 changed files with 262 additions and 8 deletions
@@ -19,18 +19,22 @@ THE SOFTWARE.

 /**
 Testcase Scenarios :
- 1) Add multiple Memcpy nodes to graph and verify node execution is
- working as expected.
+1) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D).
+Memcpy nodes are added and assigned to default device.
+2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and verify the results.
+3) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D).
+Memcpy nodes are added and assigned to Peer device.
+4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and verify the results.
 */

 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>

-/**
- * Functional Test adds memcpy nodes of types H2D, D2D and D2H to graph
- * and verifies execution sequence by launching graph.
- */
-TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
+#define ZSIZE 32
+#define YSIZE 32
+#define XSIZE 32
+
+void validateMemcpyNode3DArray(bool peerAccess = false) {
  constexpr int width{10}, height{10}, depth{10};
  hipArray *devArray1, *devArray2;
  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
@@ -42,6 +46,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
  hipStream_t streamForGraph;
  hipGraphExec_t graphExec;

+  HIP_CHECK(hipSetDevice(0));
  int *hData = reinterpret_cast<int*>(malloc(size));
  int *hOutputData = reinterpret_cast<int *>(malloc(size));

@@ -69,6 +74,12 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
                       make_hipExtent(width, height, depth), hipArrayDefault));
  HIP_CHECK(hipGraphCreate(&graph, 0));

+  // For peer access test, Memory is allocated on device(0)
+  // while memcpy nodes are allocated and assigned to peer device(1)
+  if (peerAccess) {
+    HIP_CHECK(hipSetDevice(1));
+  }
+
  // Host to Device
  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
  myparams.srcPos = make_hipPos(0, 0, 0);
@@ -79,7 +90,6 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
  myparams.dstArray = devArray1;
  myparams.kind = hipMemcpyHostToDevice;

-
  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
  dependencies.push_back(memcpyNode);

@@ -126,3 +136,247 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") {
  free(hData);
  free(hOutputData);
 }
+
+void validateMemcpyNode2DArray(bool peerAccess = false) {
+  int harray2D[YSIZE][XSIZE]{};
+  int harray2Dres[YSIZE][XSIZE]{};
+  constexpr int width{XSIZE}, height{YSIZE};
+  hipArray *devArray1, *devArray2;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  std::vector<hipGraphNode_t> dependencies;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipSetDevice(0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  // Initialize 2D object
+  for (int i = 0; i < YSIZE; i++) {
+    for (int j = 0; j < XSIZE; j++) {
+      harray2D[i][j] = i + j + 1;
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                                          0, 0, 0, formatKind);
+  // Allocate 2D device array by passing depth(0)
+  HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
+                       make_hipExtent(width, height, 0), hipArrayDefault));
+  HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc,
+                       make_hipExtent(width, height, 0), hipArrayDefault));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // For peer access test, Memory is allocated on device(0)
+  // while memcpy nodes are allocated and assigned to peer device(1)
+  if (peerAccess) {
+    HIP_CHECK(hipSetDevice(1));
+  }
+
+  // Host to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, height, 1);
+  myparams.srcPtr = make_hipPitchedPtr(harray2D, width * sizeof(int),
+                                      width, height);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
+  dependencies.push_back(memcpyNode);
+
+  // Device to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.srcArray = devArray1;
+  myparams.dstArray = devArray2;
+  myparams.extent = make_hipExtent(width, height, 1);
+  myparams.kind = hipMemcpyDeviceToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                             dependencies.size(), &myparams));
+  dependencies.clear();
+  dependencies.push_back(memcpyNode);
+
+  // Device to host
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, height, 1);
+  myparams.dstPtr = make_hipPitchedPtr(harray2Dres, width * sizeof(int),
+                                      width, height);
+  myparams.srcArray = devArray2;
+  myparams.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                             dependencies.size(), &myparams));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Validate result
+  for (int i = 0; i < YSIZE; i++) {
+    for (int j = 0; j < XSIZE; j++) {
+      if (harray2D[i][j] != harray2Dres[i][j]) {
+        INFO("harray2D: " << harray2D[i][j] << "harray2Dres: "
+              << harray2Dres[i][j] << " mismatch at (i,j) : " << i << j);
+        REQUIRE(false);
+      }
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  hipFreeArray(devArray1);
+  hipFreeArray(devArray2);
+}
+
+void validateMemcpyNode1DArray(bool peerAccess = false) {
+  int harray1D[XSIZE]{};
+  int harray1Dres[XSIZE]{};
+  constexpr int width{XSIZE};
+  hipArray *devArray1, *devArray2;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  hipGraph_t graph;
+  hipGraphNode_t memcpyNode;
+  std::vector<hipGraphNode_t> dependencies;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+
+  HIP_CHECK(hipSetDevice(0));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  // Initialize 1D object
+  for (int i = 0; i < XSIZE; i++) {
+    harray1D[i] = i + 1;
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                                          0, 0, 0, formatKind);
+  // Allocate 1D device array by passing depth(0), height(0)
+  HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+  HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc,
+                       make_hipExtent(width, 0, 0), hipArrayDefault));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // For peer access test, Memory is allocated on device(0)
+  // while memcpy nodes are allocated and assigned to peer device(1)
+  if (peerAccess) {
+    HIP_CHECK(hipSetDevice(1));
+  }
+
+  // Host to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.srcPtr = make_hipPitchedPtr(harray1D, width * sizeof(int),
+                                      width, 1);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams));
+  dependencies.push_back(memcpyNode);
+
+  // Device to Device
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.srcArray = devArray1;
+  myparams.dstArray = devArray2;
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.kind = hipMemcpyDeviceToDevice;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                             dependencies.size(), &myparams));
+  dependencies.clear();
+  dependencies.push_back(memcpyNode);
+
+  // Device to host
+  memset(&myparams, 0x0, sizeof(hipMemcpy3DParms));
+  myparams.srcPos = make_hipPos(0, 0, 0);
+  myparams.dstPos = make_hipPos(0, 0, 0);
+  myparams.extent = make_hipExtent(width, 1, 1);
+  myparams.dstPtr = make_hipPitchedPtr(harray1Dres, width * sizeof(int),
+                                      width, 1);
+  myparams.srcArray = devArray2;
+  myparams.kind = hipMemcpyDeviceToHost;
+
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(),
+                                              dependencies.size(), &myparams));
+
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Validate result
+  for (int i = 0; i < XSIZE; i++) {
+    if (harray1D[i] != harray1Dres[i]) {
+      INFO("harray1D: " << harray1D[i] << " harray1Dres: " << harray1Dres[i]
+            << " mismatch at : " << i);
+      REQUIRE(false);
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  hipFreeArray(devArray1);
+  hipFreeArray(devArray2);
+}
+
+/**
+ * Basic Functional Tests adds memcpy nodes of types H2D, D2D and D2H to graph
+ * and verifies execution sequence by launching graph on default device.
+ * Tests also verify memcpy node addition with 1D, 2D and 3D objects.
+ */
+TEST_CASE("Unit_hipGraphAddMemcpyNode_BasicFunctional") {
+  SECTION("Memcpy with 3D array on default device") {
+    validateMemcpyNode3DArray();
+  }
+
+  SECTION("Memcpy with 2D array on default device") {
+    validateMemcpyNode2DArray();
+  }
+
+  SECTION("Memcpy with 1D array on default device") {
+    validateMemcpyNode1DArray();
+  }
+}
+
+/**
+ * Peer access tests adds and assigns memcpy nodes of types H2D, D2D and D2H
+ * to peer device. Memory allocations happen on device(0) and memcpy operations
+ * are performed from device(1).
+ * Tests also verify memcpy node addition with 1D, 2D and 3D objects.
+ */
+TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") {
+  int numDevices{}, peerAccess{};
+  HIP_CHECK(hipGetDeviceCount(&numDevices));
+  if (numDevices > 1) {
+    HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0));
+  }
+
+  if (!peerAccess) {
+    WARN("Skipping test as peer device access is not found!");
+    return;
+  }
+
+  SECTION("Memcpy with 3D array on peer device") {
+    validateMemcpyNode3DArray(true);
+  }
+
+  SECTION("Memcpy with 2D array on peer device") {
+    validateMemcpyNode2DArray(true);
+  }
+
+  SECTION("Memcpy with 1D array on peer device") {
+    validateMemcpyNode1DArray(true);
+  }
+}