From 978b02c52f39f05393f6534334e530ac034eac97 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 19 Jul 2022 19:27:29 +0530 Subject: [PATCH] SWDEV-306122 - [catch2][dtest] API hipGraphKernelNodeGetParams, hipGraphKernelNodeSetParams, hipGraphExecKernelNodeSetParams. (#2802) Change-Id: Ie5eb8949f589ac7ac5eb0f02e99f2012d9c0b926 --- .../config/config_amd_windows.json | 5 +- tests/catch/include/hip_test_checkers.hh | 7 + tests/catch/include/hip_test_kernels.hh | 8 + tests/catch/unit/graph/CMakeLists.txt | 3 + .../graph/hipGraphExecKernelNodeSetParams.cc | 186 ++++++++++++++++++ .../unit/graph/hipGraphKernelNodeGetParams.cc | 159 +++++++++++++++ .../unit/graph/hipGraphKernelNodeSetParams.cc | 145 ++++++++++++++ 7 files changed, 512 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc create mode 100644 tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc create mode 100644 tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 547d92a4bf..06aa6bd2db 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -46,6 +46,9 @@ "Unit_hipArrayCreate_happy - float", "Unit_hipArrayCreate_happy - float2", "Unit_hipArrayCreate_happy - float4", - "Unit_hipMemVmm_Basic" + "Unit_hipMemVmm_Basic", + "Unit_hipGraphKernelNodeSetParams_Functional", + "Unit_hipGraphExecKernelNodeSetParams_Negative", + "Unit_hipGraphExecKernelNodeSetParams_Functional" ] } diff --git a/tests/catch/include/hip_test_checkers.hh b/tests/catch/include/hip_test_checkers.hh index fce61fe677..3e152f099b 100644 --- a/tests/catch/include/hip_test_checkers.hh +++ b/tests/catch/include/hip_test_checkers.hh @@ -98,6 +98,13 @@ size_t checkVectorADD(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch = A_h, B_h, result_H, N, [](T a, T b) { return a + b; }, expectMatch, reportMismatch); } +template +size_t checkVectorSUB(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch = true, + bool reportMismatch = true) { + return checkVectors( + A_h, B_h, result_H, N, [](T a, T b) { return a - b; }, expectMatch, reportMismatch); +} + template void checkTest(T* expected_H, T* result_H, size_t N, bool expectMatch = true) { checkVectors( diff --git a/tests/catch/include/hip_test_kernels.hh b/tests/catch/include/hip_test_kernels.hh index d59fd41b5a..68f452ba53 100644 --- a/tests/catch/include/hip_test_kernels.hh +++ b/tests/catch/include/hip_test_kernels.hh @@ -34,6 +34,14 @@ template __global__ void vectorADD(const T* A_d, const T* B_d, T* C } } +template __global__ void vectorSUB(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < NELEM; i += stride) { + C_d[i] = A_d[i] - B_d[i]; + } +} template __global__ void vectorADDReverse(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 33a6ce3cd5..a3c96eccad 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -68,6 +68,9 @@ set(TEST_SRC hipGraphAddKernelNode.cc hipGraphMemcpyNodeGetParams.cc hipGraphMemcpyNodeSetParams.cc + hipGraphKernelNodeGetParams.cc + hipGraphKernelNodeSetParams.cc + hipGraphExecKernelNodeSetParams.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc b/tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc new file mode 100644 index 0000000000..e1141fbe2a --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc @@ -0,0 +1,186 @@ +/* +Copyright (c) 2022 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 +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass hGraphExec as nullptr and verify api returns error code. +2) Pass node as nullptr and verify api returns error code. +3) Pass NodeParams as un-initialized structure object and verify api returns error code. +4) Pass pNodeParams as nullptr and verify api returns error code. +5) Pass NodeParams:func datamember as nullptr and verify api returns error code. +Functional - +1) Instantiate a graph with kernel node, obtain executable graph and update + the kernel node params with set and check it is taking effect. +*/ + +#include +#include +#include + +/** + * Negative Test for API hipGraphExecKernelNodeSetParams + */ +TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyNode, kNode{}; + hipKernelNodeParams kNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + + SECTION("Pass hipGraphExec as nullptr") { + ret = hipGraphExecKernelNodeSetParams(nullptr, kNode, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Node as nullptr") { + ret = hipGraphExecKernelNodeSetParams(graphExec, nullptr, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_AMD + /* NodeParams null check is disabled on Nvedia as + * this call gives SIGSEGV error in CUDA setup */ + SECTION("Pass NodeParams as nullptr") { + ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif +/* For below 2 scenarios - + In AMD setup this API return - hipErrorInvalidValue and + In CUDA setup this API return - hipErrorInvalidDeviceFunction + As per Cuda spec API can only return "cudaSuccess, cudaErrorInvalidValue". +*/ + SECTION("Pass NodeParams as un-initialized structure object") { + hipKernelNodeParams kNodeParams1{}; + ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams1); + REQUIRE(hipSuccess != ret); + } + SECTION("Pass NodeParams func datamember as nullptr") { + kNodeParams.func = nullptr; + ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams); + REQUIRE(hipSuccess != ret); + } + + 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(hipStreamDestroy(streamForGraph)); +} + + +/** + * Functional Test for API Exec Kernel Params + */ + +TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memcpyNode, kNode; + hipKernelNodeParams kNodeParams{}, kNodeParams1{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), + dependencies.size(), &kNodeParams)); + + memset(&kNodeParams1, 0, sizeof(kNodeParams1)); + kNodeParams1.func = reinterpret_cast(HipTest::vectorSUB); + kNodeParams1.gridDim = dim3(blocks); + kNodeParams1.blockDim = dim3(threadsPerBlock); + kNodeParams1.sharedMemBytes = 0; + kNodeParams1.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams1.extra = nullptr; + + dependencies.clear(); + dependencies.push_back(kNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + REQUIRE(hipSuccess == hipGraphExecKernelNodeSetParams(graphExec, kNode, + &kNodeParams1)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorSUB(A_h, B_h, C_h, N); + + 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(hipStreamDestroy(streamForGraph)); +} diff --git a/tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc b/tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc new file mode 100644 index 0000000000..1ccdc62b59 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphKernelNodeGetParams.cc @@ -0,0 +1,159 @@ +/* +Copyright (c) 2022 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 +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass node as nullptr and verify api returns error code. +2) Pass pNodeParams as nullptr and verify api returns error code. +Functional - +1) Create a graph, add kernel node to graph with desired kernel node params. + Verify api fetches the node params mentioned while adding kernel node. +2) Set kernel node params with hipGraphKernelNodeSetParams, + now get the params and verify both are same. +*/ + +#include +#include + +#define THREADS_PER_BLOCK 512 + +/* Test verifies hipGraphKernelNodeGetParams API Negative scenarios. + */ + +TEST_CASE("Unit_hipGraphKernelNodeGetParams_Negative") { + constexpr int N = 1024; + size_t NElem{N}; + int *A_d, *B_d, *C_d; + hipError_t ret; + hipGraph_t graph; + hipGraphNode_t kNode; + hipKernelNodeParams kNodeParams{}; + HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1); + kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); + + SECTION("Pass node as nullptr") { + ret = hipGraphKernelNodeGetParams(nullptr, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + + SECTION("Pass kNodeParams as nullptr") { + ret = hipGraphKernelNodeGetParams(kNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +static bool dim3_compare(dim3 node1, dim3 node2) { + if ((node1.x == node2.x) && (node1.y == node2.y) && (node1.z == node2.z)) + return true; + else + return false; +} + +static bool kernelParam_compare(void **p1, void ** p2) { + for (int i = 0; i < 4; i++) { + if (*reinterpret_cast(p1[i]) != *reinterpret_cast(p2[i])) + return false; + } + return true; +} + +static bool node_compare(hipKernelNodeParams *kNode1, + hipKernelNodeParams *kNode2) { + if (!dim3_compare(kNode1->blockDim, kNode2->blockDim)) + return false; + if (kNode1->extra != kNode2->extra) + return false; + if (kNode1->func != kNode2->func) + return false; + if (!dim3_compare(kNode1->gridDim, kNode2->gridDim)) + return false; + if (!kernelParam_compare(kNode1->kernelParams, kNode2->kernelParams)) + return false; + if (kNode1->sharedMemBytes != kNode2->sharedMemBytes) + return false; + return true; +} + +/* Test verifies hipGraphKernelNodeGetParams API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphKernelNodeGetParams_Functional") { + constexpr int N = 1024; + size_t NElem{N}; + int *A_d, *B_d, *C_d; + hipGraph_t graph; + hipGraphNode_t kNode; + hipKernelNodeParams kNodeParams{}; + HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1); + kNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); + + SECTION("Get Kernel Param and verify.") { + hipKernelNodeParams kNodeGetParams; + HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams)); + REQUIRE(true == node_compare(&kNodeParams, &kNodeGetParams)); + } + + SECTION("Set kernel node params then Get Kernel Param and verify.") { + hipKernelNodeParams kNodeParams1; + kNodeParams1.func = + reinterpret_cast(HipTest::vectorADDReverse); + kNodeParams1.gridDim = dim3(N / THREADS_PER_BLOCK, 1, 1); + kNodeParams1.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kNodeParams1.sharedMemBytes = 0; + kNodeParams1.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams1.extra = nullptr; + + hipKernelNodeParams kNodeGetParams1; + HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1)); + HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams1)); + REQUIRE(true == node_compare(&kNodeParams1, &kNodeGetParams1)); + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc b/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc new file mode 100644 index 0000000000..05a6ec345f --- /dev/null +++ b/tests/catch/unit/graph/hipGraphKernelNodeSetParams.cc @@ -0,0 +1,145 @@ +/* +Copyright (c) 2022 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 +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass node as nullptr and verify api returns error code. +2) Pass pNodeParams as nullptr and verify api returns error code. +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. +*/ + +#include +#include +#include + +/* Test verifies hipGraphKernelNodeSetParams API Negative scenarios. + */ + +TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") { + constexpr int N = 1024; + size_t NElem{N}; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + int *A_d, *B_d, *C_d; + hipError_t ret; + hipGraph_t graph; + hipGraphNode_t kNode; + hipKernelNodeParams kNodeParams{}; + HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); + + SECTION("Pass node as nullptr") { + ret = hipGraphKernelNodeSetParams(nullptr, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + + SECTION("Pass kNodeParams as nullptr") { + ret = hipGraphKernelNodeSetParams(kNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test for API Set Kernel Params + */ + +TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memcpyNode, kNode; + hipKernelNodeParams kNodeParams{}, kNodeParams1{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), + dependencies.size(), &kNodeParams)); + + kNodeParams1.func = reinterpret_cast(HipTest::vectorSUB); + kNodeParams1.gridDim = dim3(blocks); + kNodeParams1.blockDim = dim3(threadsPerBlock); + kNodeParams1.sharedMemBytes = 0; + kNodeParams1.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1)); + + dependencies.clear(); + dependencies.push_back(kNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorSUB(A_h, B_h, C_h, N); + + 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(hipStreamDestroy(streamForGraph)); +}