From 22fc2dd74bf3d6f299c252a454bd178234565928 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 24 Mar 2023 15:59:08 +0530 Subject: [PATCH] SWDEV-357695 - [catch2][dtest] hipGraphKernelNodeCopyAttributes API test (#90) Change-Id: I8a7494cbb056df8b2b6b0bf5810f67e4eed4db7b [ROCm/hip-tests commit: 2cf38e9e33afa8caa2bd33d09273ad78852e3332] --- .../hip-tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphKernelNodeCopyAttributes.cc | 165 ++++++++++++++++++ 2 files changed, 166 insertions(+) create mode 100644 projects/hip-tests/catch/unit/graph/hipGraphKernelNodeCopyAttributes.cc diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index ee58c1672c..5a93f0c83e 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -104,6 +104,7 @@ set(TEST_SRC hipGraphDestroy.cc hipGraphExecDestroy.cc hipGraphUpload.cc + hipGraphKernelNodeCopyAttributes.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/projects/hip-tests/catch/unit/graph/hipGraphKernelNodeCopyAttributes.cc b/projects/hip-tests/catch/unit/graph/hipGraphKernelNodeCopyAttributes.cc new file mode 100644 index 0000000000..7ebd9953ff --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipGraphKernelNodeCopyAttributes.cc @@ -0,0 +1,165 @@ +/* +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. +*/ + +#include +#include +#include + +/** + * Functional Test for API - hipGraphKernelNodeCopyAttributes + - Create graph with 2 kernel node. + Copy the attribute from 1st kernel node to 2nd kernel node. + */ + +TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_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{}; + hipStream_t stream; + 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(&stream)); + 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)); + dependencies.clear(); + dependencies.push_back(kNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + hipGraphNode_t kNode2; + + HIP_CHECK(hipGraphAddKernelNode(&kNode2, graph, nullptr, 0, &kNodeParams)); + HIP_CHECK(hipGraphKernelNodeCopyAttributes(kNode2, kNode)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Verify graph execution result + HipTest::checkVectorADD(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(stream)); +} + +/** + Negative Test for API - hipGraphKernelNodeCopyAttributes + 1) Pass source kernel node as nullptr for copy attribute api + 2) Pass destination kernel node as nullptr for copy attribute api + 3) Pass source kernel node as Uninitialize for copy attribute api + 4) Pass dest kernel node as Uninitialize for copy attribute api + */ + +TEST_CASE("Unit_hipGraphKernelNodeCopyAttributes_Attribute_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; + hipGraphNode_t memcpyNode, kNode, kNode_2; + hipKernelNodeParams kNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies; + size_t NElem{N}; + hipError_t ret; + + 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)); + dependencies.clear(); + dependencies.push_back(kNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddKernelNode(&kNode_2, graph, nullptr, 0, &kNodeParams)); + + SECTION("Pass source kernel node as nullptr for copy attribute api") { + ret = hipGraphKernelNodeCopyAttributes(nullptr, kNode); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass destination kernel node as nullptr for copy attribute api") { + ret = hipGraphKernelNodeCopyAttributes(kNode_2, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass source kernel node as Uninitialize for copy attribute api") { + hipGraphNode_t kNodeUninit{}; + ret = hipGraphKernelNodeCopyAttributes(kNodeUninit, kNode); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass dest kernel node as Uninitialize for copy attribute api") { + hipGraphNode_t kNodeUninit{}; + ret = hipGraphKernelNodeCopyAttributes(kNode_2, kNodeUninit); + REQUIRE(hipErrorInvalidValue == ret); + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} +