From 3cf8079c4cfe2bc6cce7365a605def3976c8d136 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 1 Apr 2022 05:36:26 +0530 Subject: [PATCH] SWDEV-306122 - [catch2][dtest] hipGraphMemsetNodeGetParams/hipGraphMemsetNodeSetParams APIs (#2586) Change-Id: Ib45b660ad94ec3231082deb8a0b6f62b278d2006 --- tests/catch/unit/graph/CMakeLists.txt | 2 + .../unit/graph/hipGraphMemsetNodeGetParams.cc | 141 +++++++++++++ .../unit/graph/hipGraphMemsetNodeSetParams.cc | 185 ++++++++++++++++++ 3 files changed, 328 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc create mode 100644 tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 301c4cd2db..26f635f11d 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -48,6 +48,8 @@ set(TEST_SRC hipGraphInstantiate.cc hipGraphExecUpdate.cc hipGraphExecEventRecordNodeSetEvent.cc + hipGraphMemsetNodeGetParams.cc + hipGraphMemsetNodeSetParams.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc b/tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc new file mode 100644 index 0000000000..03cbcaaf2c --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemsetNodeGetParams.cc @@ -0,0 +1,141 @@ +/* +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 Memset node to graph with desired node params. + Verify api fetches the node params mentioned while adding Memset node. +2) Set Memset node params with hipGraphMemsetNodeSetParams, + now get the params and verify both are same. +*/ + +#include + +/* Test verifies hipGraphMemsetNodeGetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Negative") { + hipError_t ret; + hipGraph_t graph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + char *devData; + HIP_CHECK(hipMalloc(&devData, 1024)); + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = 1024; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + SECTION("Pass node as nullptr") { + ret = hipGraphMemsetNodeGetParams(nullptr, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass GetNodeParams as nullptr") { + ret = hipGraphMemsetNodeGetParams(memsetNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemsetNodeGetParams API Functional scenarios. + */ + +bool memsetNodeCompare(hipMemsetParams *mNode1, hipMemsetParams *mNode2) { + if (mNode1->dst != mNode2->dst) + return false; + if (mNode1->elementSize != mNode2->elementSize) + return false; + if (mNode1->height != mNode2->height) + return false; + if (mNode1->pitch != mNode2->pitch) + return false; + if (mNode1->value != mNode2->value) + return false; + if (mNode1->width != mNode2->width) + return false; + return true; +} + +TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + + char *devData; + HIP_CHECK(hipMalloc(&devData, Nbytes)); + + hipGraph_t graph; + hipGraphNode_t memsetNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = val; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + SECTION("Get Memset Param and verify.") { + hipMemsetParams memsetGetParams; + REQUIRE(hipSuccess == hipGraphMemsetNodeGetParams(memsetNode, + &memsetGetParams)); + // Validating the result + REQUIRE(true == memsetNodeCompare(&memsetParams, &memsetGetParams)); + } + SECTION("Set memset node params then Get and verify.") { + constexpr size_t updateVal = 2; + char *devData1; + HIP_CHECK(hipMalloc(&devData1, Nbytes)); + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData1); + memsetParams.value = updateVal; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + + hipMemsetParams memsetGetParams; + REQUIRE(hipSuccess == hipGraphMemsetNodeSetParams(memsetNode, + &memsetParams)); + REQUIRE(hipSuccess == hipGraphMemsetNodeGetParams(memsetNode, + &memsetGetParams)); + // Validating the result + REQUIRE(true == memsetNodeCompare(&memsetParams, &memsetGetParams)); + HIP_CHECK(hipFree(devData1)); + } + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} diff --git a/tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc b/tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc new file mode 100644 index 0000000000..fe188d48c8 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc @@ -0,0 +1,185 @@ +/* +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. +3) Passing hipMemsetParams::dst as nullptr should return error code. +4) Passing hipMemsetParams::element size other than 1, 2, or 4 and check. +5) Passing hipMemsetParams::height as zero and check api should return error code. +Functional - +1) Add Memset node to graph, update the node params with set and + launch the graph and check the set params are executing properly. +2) Add Memset node to graph, launch graph, then update the Memset node params + with set and launch the graph and check updated params are taking effect. +*/ + +#include +#include + +/* Test verifies hipGraphMemsetNodeSetParams API invalid params scenarios. + */ +TEST_CASE("Unit_hipGraphMemsetNodeSetParams_InvalidParams") { + hipError_t ret; + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + char *devData; + HIP_CHECK(hipMalloc(&devData, 1024)); + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = 1024; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + SECTION("Pass node as nullptr") { + ret = hipGraphMemsetNodeSetParams(nullptr, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass GetNodeParams as nullptr") { + ret = hipGraphMemsetNodeSetParams(memsetNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass dest as nullptr") { + memsetParams.dst = nullptr; + ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_NVIDIA + SECTION("Pass element size other than 1, 2, or 4") { + memsetParams.dst = reinterpret_cast(devData); + memsetParams.elementSize = 9; + ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass height as zero or negative") { + memsetParams.elementSize = 2; + memsetParams.height = 0; + ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +static void validate_result(char *hData, size_t size, char val) { + // Validating the result + for (size_t i = 0; i < size; i++) { + if (hData[i] != val) { + WARN("Validation failed at- " << i << " hData[i] " << hData[i]); + REQUIRE(false); + } + } +} + +/* Test verifies hipGraphMemsetNodeSetParams API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphMemsetNodeSetParams_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + constexpr size_t updateVal = 1; + constexpr size_t updateVal2 = 2; + char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + char *A_h{nullptr}, *B_h{nullptr}; + + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, nullptr, N, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(C_d); + memsetParams.value = val; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + SECTION("Update the memsetNode and check") { + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = updateVal; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), + dependencies.size(), &memsetParams)); + HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + validate_result(A_h, Nbytes, updateVal); + } + SECTION("Update the memsetNode again and check") { + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = updateVal2; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), + dependencies.size(), &memsetParams)); + HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost)); + validate_result(B_h, Nbytes, updateVal2); + } + + HipTest::freeArrays(A_d, B_d, C_d, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +}