From e710eeea8da9c44818bdb22dbd758a9e02c497a8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 29 Dec 2022 08:47:28 +0100 Subject: [PATCH] EXSWHTEC-169 - Implement additional tests for Kernel Graph Node APIs (#7) - Tidy up hipGraphAddKernelNode tests - Tidy up hipGraphKernelNodeGetParams tests - Tidy up hipGraphKernelNodeSetParams tests - Tidy up hipGraphExecKernelNodeSetParams tests. - Disable failing test sections on AMD. --- catch/unit/graph/hipGraphAddKernelNode.cc | 72 +++++---- .../graph/hipGraphExecKernelNodeSetParams.cc | 133 ++++++++-------- .../unit/graph/hipGraphKernelNodeGetParams.cc | 75 +++++---- .../unit/graph/hipGraphKernelNodeSetParams.cc | 143 +++++++++--------- 4 files changed, 215 insertions(+), 208 deletions(-) diff --git a/catch/unit/graph/hipGraphAddKernelNode.cc b/catch/unit/graph/hipGraphAddKernelNode.cc index 79ed3c76ea..1f4d96f292 100644 --- a/catch/unit/graph/hipGraphAddKernelNode.cc +++ b/catch/unit/graph/hipGraphAddKernelNode.cc @@ -6,8 +6,10 @@ 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 @@ -22,7 +24,6 @@ THE SOFTWARE. /* Test verifies hipGraphAddKernelNode API Negative scenarios. */ - TEST_CASE("Unit_hipGraphAddKernelNode_Negative") { constexpr int N = 1024; size_t NElem{N}; @@ -31,7 +32,6 @@ TEST_CASE("Unit_hipGraphAddKernelNode_Negative") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); int *A_d, *B_d, *C_d; hipGraph_t graph; - hipError_t ret; hipGraphNode_t kNode; hipKernelNodeParams kNodeParams{}; std::vector dependencies; @@ -41,61 +41,67 @@ TEST_CASE("Unit_hipGraphAddKernelNode_Negative") { 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); + 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; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); SECTION("Pass pGraphNode as nullptr") { - ret = hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams), + hipErrorInvalidValue); } + SECTION("Pass Graph as nullptr") { - ret = hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams), + hipErrorInvalidValue); } + SECTION("Pass invalid numDependencies") { - ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams), + hipErrorInvalidValue); } + SECTION("Pass invalid numDependencies and valid list for dependencies") { HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); dependencies.push_back(kNode); - ret = hipGraphAddKernelNode(&kNode, graph, - dependencies.data(), dependencies.size()+1, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), + dependencies.size() + 1, &kNodeParams), + hipErrorInvalidValue); } + SECTION("Pass NodeParams as nullptr") { - ret = hipGraphAddKernelNode(&kNode, graph, - dependencies.data(), dependencies.size(), nullptr); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR( + hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), nullptr), + hipErrorInvalidValue); } - SECTION("Pass NodeParams func datamember as nullptr") { + +#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue + SECTION("Pass NodeParams func data member as nullptr") { kNodeParams.func = nullptr; - ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); - REQUIRE(hipSuccess != ret); + HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams), + hipErrorInvalidDeviceFunction); } - SECTION("Pass kernelParams datamember as nullptr") { - kNodeParams.func = reinterpret_cast(HipTest::vectorADD); +#endif + + SECTION("Pass kernelParams data member as nullptr") { kNodeParams.kernelParams = nullptr; - ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams), + hipErrorInvalidValue); } -#if HT_AMD -// On Cuda setup this test case getting failed + +#if HT_AMD // On Cuda setup this test case getting failed SECTION("Try adding kernel node after destroy the already created graph") { - kNodeParams.kernelParams = reinterpret_cast(kernelArgs); - HIP_CHECK(hipGraphDestroy(graph)); - ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + hipGraph_t destroyed_graph; + HIP_CHECK(hipGraphCreate(&destroyed_graph, 0)); + HIP_CHECK(hipGraphDestroy(destroyed_graph)); + HIP_CHECK_ERROR(hipGraphAddKernelNode(&kNode, destroyed_graph, nullptr, 0, &kNodeParams), + hipErrorInvalidValue); } #endif HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipGraphDestroy(graph)); } - diff --git a/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc b/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc index e1141fbe2a..524e38b460 100644 --- a/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc +++ b/catch/unit/graph/hipGraphExecKernelNodeSetParams.cc @@ -6,25 +6,27 @@ 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 + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY 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 +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 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 +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 : +Test Case 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. +5) Pass NodeParams:func data member 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. @@ -39,12 +41,11 @@ Functional - */ 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; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); hipGraph_t graph; - hipError_t ret; - hipGraphNode_t memcpyNode, kNode{}; + hipGraphNode_t kNode{}; hipKernelNodeParams kNodeParams{}; hipStream_t streamForGraph; int *A_d, *B_d, *C_d; @@ -55,57 +56,67 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") { 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); + 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)); + + hipGraphNode_t empty_node; + HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, &kNode, 1)); // Instantiate and launch the graph - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); SECTION("Pass hipGraphExec as nullptr") { - ret = hipGraphExecKernelNodeSetParams(nullptr, kNode, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(nullptr, kNode, &kNodeParams), + hipErrorInvalidValue); } + SECTION("Pass Node as nullptr") { - ret = hipGraphExecKernelNodeSetParams(graphExec, nullptr, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, nullptr, &kNodeParams), + hipErrorInvalidValue); } + #if HT_AMD - /* NodeParams null check is disabled on Nvedia as + /* NodeParams null check is disabled on Nvidia as * this call gives SIGSEGV error in CUDA setup */ SECTION("Pass NodeParams as nullptr") { - ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, nullptr); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, nullptr), + hipErrorInvalidValue); } #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") { + +#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue + SECTION("Pass NodeParams func data member as nullptr") { kNodeParams.func = nullptr; - ret = hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams); - REQUIRE(hipSuccess != ret); + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams), + hipErrorInvalidDeviceFunction); + } +#endif + +#if HT_NVIDIA // segfaults on AMD + SECTION("Pass kernelParams data member as nullptr") { + kNodeParams.kernelParams = nullptr; + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams), + hipErrorInvalidValue); + } +#endif + +#if HT_NVIDIA // segfaults on AMD + SECTION("node is not a kernel node") { + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, empty_node, &kNodeParams), + hipErrorInvalidValue); + } +#endif + + SECTION("node is not instantiated") { + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); + HIP_CHECK_ERROR(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams), + hipErrorInvalidValue); } HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); @@ -114,16 +125,15 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Negative") { 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; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); hipGraph_t graph; hipGraphNode_t memcpyNode, kNode; hipKernelNodeParams kNodeParams{}, kNodeParams1{}; @@ -136,43 +146,36 @@ TEST_CASE("Unit_hipGraphExecKernelNodeSetParams_Functional") { 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)); + + 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)); + 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); + 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)); + HIP_CHECK( + hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams)); - memset(&kNodeParams1, 0, sizeof(kNodeParams1)); - kNodeParams1.func = reinterpret_cast(HipTest::vectorSUB); + 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)); + 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(hipGraphExecKernelNodeSetParams(graphExec, kNode, &kNodeParams1)); HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); HIP_CHECK(hipStreamSynchronize(streamForGraph)); diff --git a/catch/unit/graph/hipGraphKernelNodeGetParams.cc b/catch/unit/graph/hipGraphKernelNodeGetParams.cc index 1ccdc62b59..8105768655 100644 --- a/catch/unit/graph/hipGraphKernelNodeGetParams.cc +++ b/catch/unit/graph/hipGraphKernelNodeGetParams.cc @@ -6,8 +6,10 @@ 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 @@ -18,7 +20,7 @@ THE SOFTWARE. */ /** -Testcase Scenarios : +Test Case Scenarios : Negative - 1) Pass node as nullptr and verify api returns error code. 2) Pass pNodeParams as nullptr and verify api returns error code. @@ -36,40 +38,42 @@ Functional - /* 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); + 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; - + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); SECTION("Pass node as nullptr") { - ret = hipGraphKernelNodeGetParams(nullptr, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphKernelNodeGetParams(nullptr, &kNodeParams), hipErrorInvalidValue); } SECTION("Pass kNodeParams as nullptr") { - ret = hipGraphKernelNodeGetParams(kNode, nullptr); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphKernelNodeGetParams(kNode, nullptr), hipErrorInvalidValue); } +#if HT_NVIDIA // segfaults on AMD + SECTION("node is not a kernel node") { + hipGraphNode_t empty_node; + HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0)); + HIP_CHECK_ERROR(hipGraphKernelNodeGetParams(empty_node, &kNodeParams), hipErrorInvalidValue); + } +#endif + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); @@ -83,28 +87,20 @@ static bool dim3_compare(dim3 node1, dim3 node2) { return false; } -static bool kernelParam_compare(void **p1, void ** p2) { +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; + 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; +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; } @@ -121,37 +117,36 @@ TEST_CASE("Unit_hipGraphKernelNodeGetParams_Functional") { 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)}; + 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; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); 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)); + REQUIRE(node_compare(&kNodeParams, &kNodeGetParams)); } SECTION("Set kernel node params then Get Kernel Param and verify.") { hipKernelNodeParams kNodeParams1; - kNodeParams1.func = - reinterpret_cast(HipTest::vectorADDReverse); + 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.kernelParams = reinterpret_cast(kernelArgs); kNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1)); hipKernelNodeParams kNodeGetParams1; - HIP_CHECK(hipGraphKernelNodeSetParams(kNode, &kNodeParams1)); HIP_CHECK(hipGraphKernelNodeGetParams(kNode, &kNodeGetParams1)); - REQUIRE(true == node_compare(&kNodeParams1, &kNodeGetParams1)); + + REQUIRE(node_compare(&kNodeParams1, &kNodeGetParams1)); } + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); diff --git a/catch/unit/graph/hipGraphKernelNodeSetParams.cc b/catch/unit/graph/hipGraphKernelNodeSetParams.cc index 10f936878e..0685b81a76 100644 --- a/catch/unit/graph/hipGraphKernelNodeSetParams.cc +++ b/catch/unit/graph/hipGraphKernelNodeSetParams.cc @@ -6,19 +6,21 @@ 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 + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY 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 +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 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 +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 : +Test Case Scenarios : Negative - 1) Pass node as nullptr and verify api returns error code. 2) Pass pNodeParams as nullptr and verify api returns error code. @@ -30,13 +32,12 @@ Functional - hipGraphKernelNodeSetParams, finally check taking effect after launching graph. */ -#include #include +#include #include /* Test verifies hipGraphKernelNodeSetParams API Negative scenarios. */ - TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") { constexpr int N = 1024; size_t NElem{N}; @@ -44,35 +45,53 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") { 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); + 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; - + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); SECTION("Pass node as nullptr") { - ret = hipGraphKernelNodeSetParams(nullptr, &kNodeParams); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(nullptr, &kNodeParams), hipErrorInvalidValue); } SECTION("Pass kNodeParams as nullptr") { - ret = hipGraphKernelNodeSetParams(kNode, nullptr); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, nullptr), hipErrorInvalidValue); } +#if HT_NVIDIA // on AMD this returns hipErrorInvalidValue + SECTION("Pass NodeParams func data member as nullptr") { + kNodeParams.func = nullptr; + HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, &kNodeParams), + hipErrorInvalidDeviceFunction); + } +#endif + +#if HT_NVIDIA // segfaults on AMD + SECTION("Pass kernelParams data member as nullptr") { + kNodeParams.kernelParams = nullptr; + HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(kNode, &kNodeParams), hipErrorInvalidValue); + } +#endif + +#if HT_NVIDIA // segfaults on AMD + SECTION("node is not a kernel node") { + hipGraphNode_t empty_node; + HIP_CHECK(hipGraphAddEmptyNode(&empty_node, graph, nullptr, 0)); + HIP_CHECK_ERROR(hipGraphKernelNodeSetParams(empty_node, &kNodeParams), hipErrorInvalidValue); + } +#endif + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); @@ -82,12 +101,12 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Negative") { /** * 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; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); hipGraph_t graph; hipGraphNode_t memcpyNode, kNode; hipKernelNodeParams kNodeParams{}, kNodeParams1{}; @@ -100,39 +119,34 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") { 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)); + + 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)); + 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); + 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)); + HIP_CHECK( + hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams)); - kNodeParams1.func = reinterpret_cast(HipTest::vectorSUB); + 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)); + 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)); @@ -147,12 +161,12 @@ TEST_CASE("Unit_hipGraphKernelNodeSetParams_Functional") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -static __global__ void ker_vec_add(int *A, int *B) { +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) { +static __global__ void ker_vec_sub(int* A, int* B) { int i = threadIdx.x + blockDim.x * blockIdx.x; A[i] = A[i] - B[i]; } @@ -167,7 +181,7 @@ class GraphKernelNodeGetSetParam { const int blocks = (N / threadsPerBlock); hipGraphNode_t memcpyH2D_A1, memcpyH2D_A2, memcpyD2H_A3, vec_maths; hipGraph_t graph; - hipKernelNodeParams kerNodeParams { }; + hipKernelNodeParams kerNodeParams{}; int *A1_d, *A2_d, *A1_h, *A2_h, *A3_h; public: @@ -179,32 +193,26 @@ class GraphKernelNodeGetSetParam { HIP_CHECK(hipMalloc(&A2_d, Nbytes)); // Allocate host buffers A1_h = reinterpret_cast(malloc(Nbytes)); - REQUIRE(A1_h != NULL); + REQUIRE(A1_h != nullptr); A2_h = reinterpret_cast(malloc(Nbytes)); - REQUIRE(A2_h != NULL); + REQUIRE(A2_h != nullptr); A3_h = reinterpret_cast(malloc(Nbytes)); - REQUIRE(A3_h != NULL); + REQUIRE(A3_h != nullptr); // Create all the 3 level graphs HIP_CHECK(hipGraphCreate(&graph, 0)); - void *kernelArgs[] = { &A1_d, &A2_d }; + 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)); + 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(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)); @@ -213,20 +221,18 @@ class GraphKernelNodeGetSetParam { // 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 + A1_h[i] = (rand() % 256); // NOLINT + A2_h[i] = (rand() % 256); // NOLINT } } - hipGraph_t* getRootGraph() { - return &graph; - } + 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))); + 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; @@ -246,9 +252,7 @@ class GraphKernelNodeGetSetParam { } // Function to validate result - void validateOutData() { - HipTest::checkVectorSUB(A1_h, A2_h, A3_h, N); - } + void validateOutData() { HipTest::checkVectorSUB(A1_h, A2_h, A3_h, N); } // Destroy resources ~GraphKernelNodeGetSetParam() { @@ -263,7 +267,7 @@ class GraphKernelNodeGetSetParam { }; TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") { - hipGraph_t *graph; + hipGraph_t* graph; hipStream_t streamForGraph; hipGraphExec_t graphExec; GraphKernelNodeGetSetParam GraphKernelNodeGetSetParamObj; @@ -271,8 +275,7 @@ TEST_CASE("Unit_hipGraphKernelNodeGetSetParams_Functional") { GraphKernelNodeGetSetParamObj.updateNode(); HIP_CHECK(hipStreamCreate(&streamForGraph)); // Instantiate and launch the childgraph - HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, - nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec, (*graph), nullptr, nullptr, 0)); GraphKernelNodeGetSetParamObj.fillRandInpData(); HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); HIP_CHECK(hipStreamSynchronize(streamForGraph));