SWDEV-363038 - Add Unit_hipGraphKernelNodeGetSetParams_Functional (#3062)

Add unit test: Unit_hipGraphKernelNodeGetSetParams_Functional

Add kernel node to graph with certain kernel params, now get kernel node parameters with hipGraphKernelNodeGetParams, then update the kernel node params with hipGraphKernelNodeSetParams, finally check taking effect after launching graph.

Change-Id: I2216f72f4dade6dd37663d3b0d821cb3d35d7856

[ROCm/hip commit: 58790104f2]

projects/hip/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc

@@ -25,6 +25,9 @@ Negative -
 Functional -
 1) Add kernel node to graph with certain kernel params, now update the kernel
    node params with set and check taking effect after launching graph.
+2) Add kernel node to graph with certain kernel params, now get kernel node parameters
+   with hipGraphKernelNodeGetParams, then update the kernel node params with
+   hipGraphKernelNodeSetParams, finally check taking effect after launching graph.
 */
 
 #include <hip_test_common.hh>
@@ -143,3 +146,137 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") {
   HIP_CHECK(hipGraphDestroy(graph));
   HIP_CHECK(hipStreamDestroy(streamForGraph));
 }
+
+static __global__ void ker_vec_add(int *A, int *B) {
+  int i = threadIdx.x + blockDim.x * blockIdx.x;
+  A[i] = A[i] + B[i];
+}
+
+static __global__ void ker_vec_sub(int *A, int *B) {
+  int i = threadIdx.x + blockDim.x * blockIdx.x;
+  A[i] = A[i] - B[i];
+}
+
+/**
+ Internal class for creating nested graphs.
+ */
+class GraphKernelNodeGetSetParam {
+  const int N = 1024;
+  size_t Nbytes;
+  const int threadsPerBlock = 256;
+  const int blocks = (N / threadsPerBlock);
+  hipGraphNode_t memcpyH2D_A1, memcpyH2D_A2, memcpyD2H_A3, vec_maths;
+  hipGraph_t graph;
+  hipKernelNodeParams kerNodeParams { };
+  int *A1_d, *A2_d, *A1_h, *A2_h, *A3_h;
+
+ public:
+  // Create a nested Graph
+  GraphKernelNodeGetSetParam() {
+    Nbytes = N * sizeof(int);
+    // Allocate device buffers
+    HIP_CHECK(hipMalloc(&A1_d, Nbytes));
+    HIP_CHECK(hipMalloc(&A2_d, Nbytes));
+    // Allocate host buffers
+    A1_h = reinterpret_cast<int*>(malloc(Nbytes));
+    REQUIRE(A1_h != NULL);
+    A2_h = reinterpret_cast<int*>(malloc(Nbytes));
+    REQUIRE(A2_h != NULL);
+    A3_h = reinterpret_cast<int*>(malloc(Nbytes));
+    REQUIRE(A3_h != NULL);
+    // Create all the 3 level graphs
+    HIP_CHECK(hipGraphCreate(&graph, 0));
+    void *kernelArgs[] = { &A1_d, &A2_d };
+    kerNodeParams.func = reinterpret_cast<void*>(ker_vec_add);
+    kerNodeParams.gridDim = dim3(blocks);
+    kerNodeParams.blockDim = dim3(threadsPerBlock);
+    kerNodeParams.sharedMemBytes = 0;
+    kerNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
+    kerNodeParams.extra = nullptr;
+    HIP_CHECK(
+        hipGraphAddKernelNode(&vec_maths, graph, nullptr, 0, &kerNodeParams));
+    // Add nodes to graph
+    HIP_CHECK(
+        hipGraphAddMemcpyNode1D(&memcpyH2D_A1, graph, nullptr, 0, A1_d, A1_h,
+                                Nbytes, hipMemcpyHostToDevice));
+    HIP_CHECK(
+        hipGraphAddMemcpyNode1D(&memcpyH2D_A2, graph, nullptr, 0, A2_d, A2_h,
+                                Nbytes, hipMemcpyHostToDevice));
+    HIP_CHECK(
+        hipGraphAddMemcpyNode1D(&memcpyD2H_A3, graph, nullptr, 0, A3_h, A1_d,
+                                Nbytes, hipMemcpyDeviceToHost));
+    HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A1, &vec_maths, 1));
+    HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A2, &vec_maths, 1));
+    HIP_CHECK(hipGraphAddDependencies(graph, &vec_maths, &memcpyD2H_A3, 1));
+  }
+
+  // Fill Random Input Data
+  void fillRandInpData() {
+    for (int i = 0; i < N; i++) {
+      A1_h[i] = (rand() % 256);  //NOLINT
+      A2_h[i] = (rand() % 256);  //NOLINT
+    }
+  }
+
+  hipGraph_t* getRootGraph() {
+    return &graph;
+  }
+
+  void updateNode() {
+    size_t numNodes = 0;
+    HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
+    hipGraphNode_t *nodes = reinterpret_cast<hipGraphNode_t*>(malloc(
+        numNodes * sizeof(hipGraphNode_t)));
+    HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
+    // Get the Graph node from the embedded graph
+    size_t nodeIdx = 0;
+    for (size_t idx = 0; idx < numNodes; idx++) {
+      hipGraphNodeType nodeType;
+      HIP_CHECK(hipGraphNodeGetType(nodes[idx], &nodeType));
+      if (nodeType == hipGraphNodeTypeKernel) {
+        nodeIdx = idx;
+        break;
+      }
+    }
+    hipKernelNodeParams nodeParam;
+    HIP_CHECK(hipGraphKernelNodeGetParams(nodes[nodeIdx], &nodeParam));
+    nodeParam.func = reinterpret_cast<void*>(ker_vec_sub);
+    HIP_CHECK(hipGraphKernelNodeSetParams(nodes[nodeIdx], &nodeParam));
+    free(nodes);
+  }
+
+  // Function to validate result
+  void validateOutData() {
+    HipTest::checkVectorSUB<int>(A1_h, A2_h, A3_h, N);
+  }
+
+  // Destroy resources
+  ~GraphKernelNodeGetSetParam() {
+    // Free all allocated buffers
+    HIP_CHECK(hipFree(A2_d));
+    HIP_CHECK(hipFree(A1_d));
+    free(A3_h);
+    free(A2_h);
+    free(A1_h);
+    HIP_CHECK(hipGraphDestroy(graph));
+  }
+};
+
+TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") {
+  hipGraph_t *graph;
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  GraphKernelNodeGetSetParam GraphKernelNodeGetSetParamObj;
+  graph = GraphKernelNodeGetSetParamObj.getRootGraph();
+  GraphKernelNodeGetSetParamObj.updateNode();
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  // Instantiate and launch the childgraph
+  HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr,
+                                nullptr, 0));
+  GraphKernelNodeGetSetParamObj.fillRandInpData();
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+  GraphKernelNodeGetSetParamObj.validateOutData();
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+}