SWDEV-373476 - [catch2][dtest] hipGraphDebugDotPrint API test case added with supported type of nodes (#248)

Change-Id: Id910fabd651a28863e67fe497cd832d8b1ea7802

[ROCm/hip-tests commit: 82e874fc0b]

projects/hip-tests/catch/unit/graph/hipGraphDebugDotPrint.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
@@ -17,6 +17,14 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+/**
+* @addtogroup hipGraphDebugDotPrint hipGraphDebugDotPrint
+* @{
+* @ingroup GraphTest
+* `hipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned int flags)` -
+* Write a DOT file describing graph structure.
+*/
+
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>
 #include <hip_test_kernels.hh>
@@ -30,6 +38,13 @@ THE SOFTWARE.
 #include <algorithm>
 #include <string>
 
+__device__ int globalIn[N];
+
+static void callbackfunc(void *A_h) {
+  int *A = reinterpret_cast<int *>(A_h);
+  std::iota(A, A+N, 0);
+}
+
 static void deleteFile(const char* fName) {
   if ( remove(fName) != 0 ) {
     INFO("Error in deleting file -" << fName);
@@ -42,8 +57,9 @@ static bool checkFileExists(const char* fName) {
   return (access(fName, F_OK) != -1);
 }
 
-static int countSubstr(std::string input_str, std::string substr) {
-  int count = 0;
+static unsigned countSubstr(const std::string &input_str,
+                            const std::string &substr) {
+  unsigned count = 0;
   std::string::size_type srch_pos = 0, cur_pos = 0;
   while ((cur_pos = input_str.find(substr, srch_pos)) != std::string::npos) {
     count++;
@@ -53,35 +69,65 @@ static int countSubstr(std::string input_str, std::string substr) {
 }
 
 static bool validateDotFile(const char* fName,
-                           std::map<std::string, int> *graphData) {
-  bool isTestPassed = true;
+                            const std::map<std::string, unsigned> &graphData) {
   std::ifstream infile(fName);
   std::stringstream buffer;
   buffer << infile.rdbuf();
-  std::map<std::string, int>::iterator it;
-  for (it = (*graphData).begin(); it != (*graphData).end(); it++) {
-    if (it->second != countSubstr(buffer.str(), it->first)) {
-      isTestPassed = false;
-      break;
+  const std::string buffer_str = buffer.str();
+  for (auto it = graphData.begin(); it != graphData.end(); it++) {
+    unsigned count = countSubstr(buffer_str, it->first);
+    if (it->second != count) {
+      INFO("validateDotFile: Failed for key :: " << it->first <<
+                   " : " << count << " Expected : " << it->second);
+      return false;
     }
   }
-  return isTestPassed;
+  return true;
 }
 
+/**
+* Test Description
+* ------------------------
+*  - Functional Test for API - hipGraphDebugDotPrint
+*   Call hipGraphDebugDotPrint and provice path where to write the DOT file.
+*   Verify that DOT file get created or not for each flag passed.
+*   1) Add MemcpyNode node to graph & validate its DebugDotPrint descriptions
+*   2) Add kernel node to graph & validate its DebugDotPrint descriptions
+*   3) Add memset node to graph & validate its DebugDotPrint descriptions
+*   4) Add emptyNode to graph & validate its DebugDotPrint descriptions
+*   5) Add childGraphNode to graph & validate its DebugDotPrint descriptions
+*   6) Add eventRecord to graph & validate its DebugDotPrint descriptions
+*   7) Add eventWait to graph & validate its DebugDotPrint descriptions
+*   8) Add hostNode to graph & validate its DebugDotPrint descriptions
+*   9) Add mecpyNode1D to graph & validate its DebugDotPrint descriptions
+*   10) Add mecpyNode3D to graph & validate its DebugDotPrint descriptions
+*   11) Add MemcpyNodeToSymbol to graph & validate its DebugDotPrint descriptions
+*   12) Add MemcpyNodeFromSymbol to graph & validate its DebugDotPrint descriptions
+*   13) Add Dependencies to graph & validate its DebugDotPrint descriptions
+* Test source
+* ------------------------
+*  - /unit/graph/hipGraphDebugDotPrint.cc
+* Test requirements
+* ------------------------
+*  - HIP_VERSION >= 5.6
+*/
+
 static void hipGraphDebugDotPrint_Functional(const char* fName,
                                              unsigned int flag = 0) {
   constexpr size_t Nbytes = N * sizeof(int);
   constexpr auto blocksPerCU = 6;  // to hide latency
   constexpr auto threadsPerBlock = 256;
-  hipGraph_t graph;
+  hipGraph_t graph, childGraph;
   hipStream_t stream;
   hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kNodeAdd, memsetNode;
+  hipGraphNode_t emptyNode, childGraphNode, eventWait, eventRecord, hostNode;
   hipKernelNodeParams kNodeParams{};
   int *A_d, *B_d, *C_d, *mem_d;
-  int *A_h, *B_h, *C_h;
+  int *A_h, *B_h, *C_h, *mem_h;
   hipGraphExec_t graphExec;
   size_t NElem{N};
 
+  mem_h = reinterpret_cast<int*>(malloc(Nbytes));
   HIP_CHECK(hipMalloc(&mem_d, Nbytes));
   HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
   unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
@@ -100,15 +146,88 @@ static void hipGraphDebugDotPrint_Functional(const char* fName,
   kNodeParams.sharedMemBytes = 0;
   kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
   kNodeParams.extra = nullptr;
+
+  // Add Kernel node to graph & validate its DebugDotPrint descriptions
   HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams));
 
+  // Add MemCpy node to graph & validate its DebugDotPrint descriptions
   HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
                                     Nbytes, hipMemcpyDeviceToHost));
 
+  // Add Dependencies to graph & validate its DebugDotPrint descriptions
   HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kNodeAdd, 1));
   HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kNodeAdd, 1));
   HIP_CHECK(hipGraphAddDependencies(graph, &kNodeAdd, &memcpy_C, 1));
 
+  // Add emptyNode to graph & validate its DebugDotPrint descriptions
+  HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, NULL, 0));
+
+  // Add hostNode to graph & validate its DebugDotPrint descriptions
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = mem_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams));
+
+  hipEvent_t event;
+  HIP_CHECK(hipEventCreate(&event));
+
+  // Add eventRecord to graph & validate its DebugDotPrint descriptions
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventRecord, graph, nullptr,
+                                       0, event));
+
+  // Add eventWait to graph & validate its DebugDotPrint descriptions
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventWait, graph, nullptr, 0, event));
+
+  HIP_CHECK(hipGraphCreate(&childGraph, 0));
+
+  // Add emcpyNode3D to graph & validate its DebugDotPrint descriptions
+  constexpr int width{10}, height{10}, depth{10};
+  hipArray *devArray1;
+  hipChannelFormatKind formatKind = hipChannelFormatKindSigned;
+  hipMemcpy3DParms myparams;
+  uint32_t size = width * height * depth * sizeof(int);
+  hipGraphNode_t mcpyNode3D;
+  int *hData = reinterpret_cast<int*>(malloc(size));
+  REQUIRE(hData != nullptr);
+
+  // Initialize host buffer
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        hData[i*width*height + j*width + k] = i*width*height + j*width + k;
+      }
+    }
+  }
+
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8,
+                                                          0, 0, 0, formatKind);
+  HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc,
+                       make_hipExtent(width, height, depth), hipArrayDefault));
+
+  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, depth);
+  myparams.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int),
+                                       width, height);
+  myparams.dstArray = devArray1;
+  myparams.kind = hipMemcpyHostToDevice;
+  HIP_CHECK(hipGraphAddMemcpyNode(&mcpyNode3D, childGraph,
+                                  nullptr, 0, &myparams));
+
+  // Add MemcpyNodeToSymbol to graph & validate its DebugDotPrint description
+  hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode;
+
+  HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, childGraph,
+                                   nullptr, 0, HIP_SYMBOL(globalIn),
+                                   B_h, Nbytes, 0, hipMemcpyHostToDevice));
+
+  // Add MemcpyNodeFromSymbol to graph & validate its DebugDotPrint description
+  HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, childGraph,
+                                     nullptr, 0, B_h, HIP_SYMBOL(globalIn),
+                                     Nbytes, 0, hipMemcpyDeviceToHost));
+
+  // Add memset node to graph & validate its DebugDotPrint descriptions
   hipMemsetParams memsetParams{};
   memset(&memsetParams, 0, sizeof(memsetParams));
   memsetParams.dst = reinterpret_cast<void*>(mem_d);
@@ -117,20 +236,42 @@ static void hipGraphDebugDotPrint_Functional(const char* fName,
   memsetParams.elementSize = sizeof(char);
   memsetParams.width = Nbytes;
   memsetParams.height = 1;
-  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, childGraph, nullptr, 0,
                                   &memsetParams));
 
-  std::map<std::string, int> graphData;
-  graphData["->"] = 3;     //  number of edges
-  graphData["MEMCPY"] = 3;
-  graphData["HtoD"] = 2;
-  graphData["DtoH"] = 1;
+  // Add childGraphNode to graph & validate its DebugDotPrint descriptions
+  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph,
+                                      nullptr, 0, childGraph));
+
+  std::map<std::string, unsigned> graphData;
+  graphData["->"] = 3;       //  number of edges
+  graphData["MEMCPY"] = 6;
+  graphData["HtoA"] = 1;
+  graphData["HtoD"] = 3;
+  graphData["DtoH"] = 2;
   graphData["MEMSET"] = 1;
-  if ( flag == hipGraphDebugDotFlagsVerbose ) graphData["KERNEL"] = 1;
+  graphData["EMPTY"] = 1;
+  graphData["EVENT_WAIT"] = 1;
+  graphData["EVENT_RECORD"] = 1;
+  graphData["subgraph"] = 2;
+  graphData["HOST"] = 1;
+
+#if HT_NVIDIA
+  if ( flag == hipGraphDebugDotFlagsVerbose ||
+       flag == hipGraphDebugDotFlagsMemcpyNodeParams ) {
+    graphData["HOST"] = 7;
+  }
+#endif
+
+  if ( flag == hipGraphDebugDotFlagsVerbose ||
+       flag == hipGraphDebugDotFlagsKernelNodeAttributes
+  ) {
+    graphData["KERNEL"] = 1;
+  }
 
   HIP_CHECK(hipGraphDebugDotPrint(graph, fName, flag));
   REQUIRE(true == checkFileExists(fName));
-  REQUIRE(true == validateDotFile(fName, &graphData));
+  REQUIRE(true == validateDotFile(fName, graphData));
   deleteFile(fName);
 
   HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
@@ -140,10 +281,12 @@ static void hipGraphDebugDotPrint_Functional(const char* fName,
   // Verify graph execution result
   HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);
 
+  free(mem_h);
   HIP_CHECK(hipFree(mem_d));
   HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
   HIP_CHECK(hipGraphExecDestroy(graphExec));
   HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphDestroy(childGraph));
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
@@ -189,17 +332,24 @@ TEST_CASE("Unit_hipGraphDebugDotPrint_Functional") {
                           hipGraphDebugDotFlagsKernelNodeAttributes);
   }
 }
-#endif  //  __linux__
 
 /**
-* Negative Test for API - hipGraphDebugDotPrint Argument Check
-1) Pass graph as nullptr.
-2) Pass graph as uninitialize structure
-3) Pass path for dot file to store as nullptr
-4) Pass path for dot file to store as empth path
-5) Pass flag as hipGraphDebugDotFlags MIN - 1
-6) Pass flag as hipGraphDebugDotFlags MAX + 1
-7) Pass flag as INT_MAX
+* Test Description
+* ------------------------
+*  - Negative Test for API - hipGraphDebugDotPrint Argument Check
+*   1) Pass graph as nullptr
+*   2) Pass graph as uninitialize structure
+*   3) Pass path for dot file to store as nullptr
+*   4) Pass path for dot file to store as empth path
+*   5) Pass flag as hipGraphDebugDotFlags MIN - 1
+*   6) Pass flag as hipGraphDebugDotFlags MAX + 1
+*   7) Pass flag as INT_MAX
+* Test source
+* ------------------------
+*  - /unit/graph/hipGraphDebugDotPrint.cc
+* Test requirements
+* ------------------------
+*  - HIP_VERSION >= 5.6
 */
 
 #define DOT_FILE_PATH_NEG   "./graphDotFileNeg.dot"
@@ -241,6 +391,7 @@ TEST_CASE("Unit_hipGraphDebugDotPrint_Argument_Check") {
     ret = hipGraphDebugDotPrint(graph, DOT_FILE_PATH_NEG, INT_MAX);
     REQUIRE(hipSuccess == ret);
   }
+  deleteFile(DOT_FILE_PATH_NEG);
   HIP_CHECK(hipGraphDestroy(graph));
 }
-
+#endif  //  __linux__