From cbf80dda940772752b84c14b8a11742a717fb938 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 21 Jun 2023 11:52:37 +0530 Subject: [PATCH] SWDEV-366390 - [catch2][dtest] Adding additional functional tests for hipGraphAddEmptyNode() (#106) Change-Id: Ibc2f72280a0c4f50c7d28acd5e5de4b6e835a3d6 --- catch/unit/graph/hipGraphAddEmptyNode.cc | 159 ++++++++++++++++++++++- 1 file changed, 156 insertions(+), 3 deletions(-) diff --git a/catch/unit/graph/hipGraphAddEmptyNode.cc b/catch/unit/graph/hipGraphAddEmptyNode.cc index b21315dd2a..23d477f595 100644 --- a/catch/unit/graph/hipGraphAddEmptyNode.cc +++ b/catch/unit/graph/hipGraphAddEmptyNode.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -21,10 +21,18 @@ THE SOFTWARE. Testcase Scenarios : 1) Create and add empty node to graph and verify addition is successful. 2) Negative Scenarios + 3) Functional Test to use empty node as barrier to wait for multiple nodes. + In a graph add 3 independent memcpy_h2d nodes, add an empty node with + dependencies on the 3 memcpy_h2d nodes, add 3 independent kernel nodes, + add another empty node with dependencies on the 3 kernel nodes and + add 3 independent memcpy_d2h nodes with dependencies on previous empty + node. Execute the graph and validate the results. */ - +#include #include +#include +#define TEST_LOOP_SIZE 50 /** * Functional Test to add empty node with dependencies */ @@ -90,7 +98,7 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") { dependencies.data(), dependencies.size())); } // pDependencies is nullptr - SECTION("Null Graph") { + SECTION("Dependencies is null") { REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(&emptyNode, graph, nullptr, dependencies.size())); } @@ -98,3 +106,148 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") { HIP_CHECK(hipFree(pOutBuff_d)); HIP_CHECK(hipGraphDestroy(graph)); } + +// Function to fill input data +static void fillRandInpData(int *A1_h, int *A2_h, int *A3_h, size_t N) { + unsigned int seed = time(nullptr); + for (size_t i = 0; i < N; i++) { + A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + A3_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + } +} + +// Function to validate result +static void validateOutData(int *A1_h, int *A2_h, size_t N) { + for (size_t i = 0; i < N; i++) { + int result = (A1_h[i]*A1_h[i]); + REQUIRE(result == A2_h[i]); + } +} +/** + * Functional Test to use empty node as barrier to wait for multiple nodes. + */ +TEST_CASE("Unit_hipGraphAddEmptyNode_BarrierFunc") { + size_t size = 1024; + constexpr auto blocksPerCU = 6; + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, size); + hipGraph_t graph; + std::vector nodeDependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr}, + *outputVec_d1{nullptr}; + int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr}, + *outputVec_d2{nullptr}; + int *inputVec_d3{nullptr}, *inputVec_h3{nullptr}, *outputVec_h3{nullptr}, + *outputVec_d3{nullptr}; + // host and device allocation + HipTest::initArrays(&inputVec_d1, &outputVec_d1, nullptr, + &inputVec_h1, &outputVec_h1, nullptr, size, false); + HipTest::initArrays(&inputVec_d2, &outputVec_d2, nullptr, + &inputVec_h2, &outputVec_h2, nullptr, size, false); + HipTest::initArrays(&inputVec_d3, &outputVec_d3, nullptr, + &inputVec_h3, &outputVec_h3, nullptr, size, false); + // add nodes to graph + hipGraphNode_t memcpyH2D_1, memcpyH2D_2, memcpyH2D_3; + hipGraphNode_t vecSqr1, vecSqr2, vecSqr3; + hipGraphNode_t memcpyD2H_1, memcpyD2H_2, memcpyD2H_3; + hipGraphNode_t emptyNode1, emptyNode2; + // Create memcpy h2d nodes + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_1, graph, nullptr, + 0, inputVec_d1, inputVec_h1, (sizeof(int)*size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_2, graph, nullptr, + 0, inputVec_d2, inputVec_h2, (sizeof(int)*size), hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_3, graph, nullptr, + 0, inputVec_d3, inputVec_h3, (sizeof(int)*size), hipMemcpyHostToDevice)); + // Create dependency list + nodeDependencies.push_back(memcpyH2D_1); + nodeDependencies.push_back(memcpyH2D_2); + nodeDependencies.push_back(memcpyH2D_3); + // Create emptyNode and add it to graph with dependency + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode1, graph, nodeDependencies.data(), + nodeDependencies.size())); + nodeDependencies.clear(); + nodeDependencies.push_back(emptyNode1); + // Creating kernel nodes + hipKernelNodeParams kerNodeParams1{}, kerNodeParams2{}, kerNodeParams3{}; + void* kernelArgs1[] = {reinterpret_cast(&inputVec_d1), + reinterpret_cast(&outputVec_d1), + reinterpret_cast(&size)}; + kerNodeParams1.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams1.gridDim = dim3(blocks); + kerNodeParams1.blockDim = dim3(threadsPerBlock); + kerNodeParams1.sharedMemBytes = 0; + kerNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); + kerNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vecSqr1, graph, nodeDependencies.data(), + nodeDependencies.size(), &kerNodeParams1)); + void* kernelArgs2[] = {reinterpret_cast(&inputVec_d2), + reinterpret_cast(&outputVec_d2), + reinterpret_cast(&size)}; + kerNodeParams2.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams2.gridDim = dim3(blocks); + kerNodeParams2.blockDim = dim3(threadsPerBlock); + kerNodeParams2.sharedMemBytes = 0; + kerNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); + kerNodeParams2.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vecSqr2, graph, nodeDependencies.data(), + nodeDependencies.size(), &kerNodeParams2)); + void* kernelArgs3[] = {reinterpret_cast(&inputVec_d3), + reinterpret_cast(&outputVec_d3), + reinterpret_cast(&size)}; + kerNodeParams3.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams3.gridDim = dim3(blocks); + kerNodeParams3.blockDim = dim3(threadsPerBlock); + kerNodeParams3.sharedMemBytes = 0; + kerNodeParams3.kernelParams = reinterpret_cast(kernelArgs3); + kerNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&vecSqr3, graph, nodeDependencies.data(), + nodeDependencies.size(), &kerNodeParams3)); + nodeDependencies.clear(); + nodeDependencies.push_back(vecSqr1); + nodeDependencies.push_back(vecSqr2); + nodeDependencies.push_back(vecSqr3); + // Create emptyNode and add it to graph with dependency + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode2, graph, nodeDependencies.data(), + nodeDependencies.size())); + nodeDependencies.clear(); + nodeDependencies.push_back(emptyNode2); + // Create memcpy d2h nodes + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph, + nodeDependencies.data(), nodeDependencies.size(), outputVec_h1, + outputVec_d1, (sizeof(int)*size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph, + nodeDependencies.data(), nodeDependencies.size(), outputVec_h2, + outputVec_d2, (sizeof(int)*size), hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_3, graph, + nodeDependencies.data(), nodeDependencies.size(), outputVec_h3, + outputVec_d3, (sizeof(int)*size), hipMemcpyDeviceToHost)); + nodeDependencies.clear(); + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec{nullptr}; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, + nullptr, 0)); + // Execute graph + for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { + fillRandInpData(inputVec_h1, inputVec_h2, inputVec_h3, size); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + validateOutData(inputVec_h1, outputVec_h1, size); + validateOutData(inputVec_h2, outputVec_h2, size); + validateOutData(inputVec_h3, outputVec_h3, size); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + // Free + HipTest::freeArrays(inputVec_d1, outputVec_d1, nullptr, + inputVec_h1, outputVec_h1, nullptr, false); + HipTest::freeArrays(inputVec_d2, outputVec_d2, nullptr, + inputVec_h2, outputVec_h2, nullptr, false); + HipTest::freeArrays(inputVec_d3, outputVec_d3, nullptr, + inputVec_h3, outputVec_h3, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +}