From 28ea47b545db366c0bab5f18ce710f4b0c0bdc50 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 25 Feb 2022 14:04:23 +0530 Subject: [PATCH] SWDEV-306122 - [catch2][dtest] Tests for hipGraphDestroyNode API (#2509) Added negative and functional tests for hipGraphDestroyNode API Change-Id: Ifc4183b94e8fd428c44f2bbc3016539f972b8069 [ROCm/hip-tests commit: 9e86513dbe23f479a7ca0afc271190a43340d5c3] --- .../hip-tests/catch/unit/graph/CMakeLists.txt | 1 + .../catch/unit/graph/hipGraphDestroyNode.cc | 139 ++++++++++++++++++ 2 files changed, 140 insertions(+) create mode 100644 projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index aa23776294..8f0857c2c9 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -35,6 +35,7 @@ set(TEST_SRC hipGraphAddMemcpyNodeToSymbol.cc hipGraphExecMemsetNodeSetParams.cc hipGraphMemcpyNodeSetParamsToSymbol.cc + hipGraphDestroyNode.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc new file mode 100644 index 0000000000..4c544a3f68 --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc @@ -0,0 +1,139 @@ +/* +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 of hipGraphDestroyNode API: + +Negative :: + +1) Pass nullptr to graph node + +Functional :: + +1) Create Node and destroy the node +2) Create graph with dependencies and destroy one of the dependency node + before executing the graph. +*/ + +#include +#include +#include + + +/* This test covers the negative scenarios of + hipGraphDestroyNode API */ +TEST_CASE("Unit_hipGraphDestroyNode_Negative") { + SECTION("Passing nullptr to graph Node") { + REQUIRE(hipGraphDestroyNode(nullptr) == hipErrorInvalidValue); + } +} + +/* This test covers the basic functionality of + hipGraphDestroyNode API where we create and destroy + the node +*/ +TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") { + char *pOutBuff_d{}; + constexpr size_t size = 1024; + hipGraph_t graph{}; + hipGraphNode_t memsetNode{}; + + HIP_CHECK(hipMalloc(&pOutBuff_d, size)); + hipMemsetParams memsetParams{}; + memsetParams.dst = reinterpret_cast(pOutBuff_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = size * sizeof(char); + memsetParams.height = 1; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + REQUIRE(hipGraphDestroyNode(memsetNode) == hipSuccess); + HIP_CHECK(hipFree(pOutBuff_d)); +} + +/* +This testcase verifies the following scenario where +graph is created with dependencies and one of the dependency is +destroyed before execute the graph +*/ +TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { + 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 memcpyH2D_A, memcpyH2D_B, memcpyH2D_B2Copies, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + hipStream_t streamForGraph; + + 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(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B2Copies, graph, nullptr, + 0, B_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, B_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B2Copies, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + // Destroy one of the dependency node + HIP_CHECK(hipGraphDestroyNode(memcpyH2D_B)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, C_h, B_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); +}