SWDEV-366390 - [catch2][dtest] Added additional functional tests for the API hipGraphLaunch() (#140)

Change-Id: I47f299731ee19428dca62d1a2d67569fdcd5bf11

[ROCm/hip-tests commit: 16e9c911ff]

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

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 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
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>
-
+#include <hip_test_kernels.hh>
 /* Test verifies hipGraphLaunch API
 Negative scenarios -
 1) Pass graphExec as nullptr and verify api returns error code.
@@ -32,8 +32,16 @@ Negative scenarios -
 Functional Scenario -
 1) Check basic functionality with stream as hipStreamPerThread
 2) Test hipGraphLaunch call on multiple devices.
+3) Create a graph with multiple nodes. Create an executable graph.
+   Launch the executable graph 3 times in stream simultaneously.
+   Wait for stream. Validate the output. No issues should be observed
+4) Create a graph with multiple nodes. Create an executable graph.
+   Verify if an executable graph be launched on null stream.
 */
 
+#define SIZE 1024
+#define TEST_LOOP_SIZE 3
+
 TEST_CASE("Unit_hipGraphLaunch_Negative") {
   hipError_t ret;
   SECTION("Pass pGraphExec as nullptr") {
@@ -305,3 +313,100 @@ TEST_CASE("Unit_hipGraphLaunch_Functional_multidevice_test") {
     SUCCEED("Skipped the testcase as there is no device to test.");
   }
 }
+
+// Function to fill input data
+static void fillRandInpData(int *A1_h, int *A2_h, size_t N) {
+  unsigned int seed = time(nullptr);
+  for (size_t i = 0; i < N; i++) {
+    A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
+    A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF);
+  }
+}
+// Function to validate result
+static void validateOutData(int *A1_h, int *A2_h, size_t N) {
+  for (size_t i = 0; i < N; i++) {
+    int result = (A1_h[i]*A1_h[i]);
+    REQUIRE(result == A2_h[i]);
+  }
+}
+/*
+ * 1.Create a graph with multiple nodes. Create an executable graph.
+ * Launch the executable graph 3 times in stream simultaneously.
+ * Wait for stream. Validate the output. No issues should be observed
+ * 2.Create a graph with multiple nodes. Create an executable graph.
+ * Verify if an executable graph be launched on null stream.
+*/
+TEST_CASE("Unit_hipGraphLaunch_Functional_MultipleLaunch") {
+  size_t memSize = SIZE;
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
+                            threadsPerBlock, SIZE);
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  int *A_h{nullptr}, *A_d{nullptr}, *C_d{nullptr}, *C_h{nullptr};
+
+  HipTest::initArrays<int>(&A_d, &C_d, nullptr,
+               &A_h, &C_h, nullptr, SIZE, false);
+
+  hipGraphNode_t memcpyH2D, memcpyD2H, kernelNode;
+
+  // Create memcpy H2D nodes
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr,
+      0, A_d, A_h, (sizeof(int)*SIZE), hipMemcpyHostToDevice));
+  nodeDependencies.push_back(memcpyH2D);
+  // Creating kernel node
+  hipKernelNodeParams kerNodeParams;
+  void* kernelArgs[] = {reinterpret_cast<void*>(&A_d),
+                        reinterpret_cast<void*>(&C_d),
+                        reinterpret_cast<void*>(&memSize)};
+  kerNodeParams.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
+  kerNodeParams.gridDim = dim3(blocks);
+  kerNodeParams.blockDim = dim3(threadsPerBlock);
+  kerNodeParams.sharedMemBytes = 0;
+  kerNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
+  kerNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(),
+                                    nodeDependencies.size(), &kerNodeParams));
+  nodeDependencies.clear();
+  nodeDependencies.push_back(kernelNode);
+
+  // Create memcpy D2H nodes
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nodeDependencies.data(),
+                         nodeDependencies.size(), C_h, C_d, (sizeof(int)*SIZE),
+                         hipMemcpyDeviceToHost));
+  nodeDependencies.clear();
+
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec{nullptr};
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  // Execute graph
+  SECTION("Multiple Graph Launch") {
+    for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) {
+      fillRandInpData(A_h, C_h, SIZE);
+      HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+      HIP_CHECK(hipStreamSynchronize(streamForGraph));
+      validateOutData(A_h, C_h, SIZE);
+    }
+  }
+  SECTION("Graph launch on Null stream") {
+    for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) {
+      fillRandInpData(A_h, C_h, SIZE);
+      HIP_CHECK(hipGraphLaunch(graphExec, 0));
+      HIP_CHECK(hipStreamSynchronize(0));
+      validateOutData(A_h, C_h, SIZE);
+    }
+  }
+
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+
+  // Free
+  HipTest::freeArrays<int>(A_d, C_d, nullptr, A_h, C_h, nullptr, false);
+}