From 9e4020871e3b4b5b3f010eaafd83f4bf9309d37f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 29 Nov 2022 07:08:27 +0530 Subject: [PATCH] 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: 58790104f20dfcd78fc832d114b40120dde9f4a9] --- .../unit/graph/hipGraphKernelNodeSetParams.cc | 137 ++++++++++++++++++ 1 file changed, 137 insertions(+) diff --git a/projects/hip/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc b/projects/hip/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc index 05a6ec345f..10f936878e 100644 --- a/projects/hip/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc +++ b/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 @@ -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(malloc(Nbytes)); + REQUIRE(A1_h != NULL); + A2_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A2_h != NULL); + A3_h = reinterpret_cast(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(ker_vec_add); + kerNodeParams.gridDim = dim3(blocks); + kerNodeParams.blockDim = dim3(threadsPerBlock); + kerNodeParams.sharedMemBytes = 0; + kerNodeParams.kernelParams = reinterpret_cast(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(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(ker_vec_sub); + HIP_CHECK(hipGraphKernelNodeSetParams(nodes[nodeIdx], &nodeParam)); + free(nodes); + } + + // Function to validate result + void validateOutData() { + HipTest::checkVectorSUB(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)); +}