From d5be1c497d4c5b0bb3fbcd55bdfadd08ead10cc5 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 17 May 2022 16:01:43 +0530 Subject: [PATCH] SWDEV-306122 - [catch2][dtest] Added following test scenarios for hipGraphEventRecordNodeSetEvent (#2634) 1. Functional Test cases 2. Negative Test cases Change-Id: Ifb851ce261a996feda1c547426f583eb3f1d57a9 --- catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphEventRecordNodeSetEvent.cc | 246 ++++++++++++++++++ 2 files changed, 247 insertions(+) create mode 100644 catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index 4be0361ff6..aa560cc8f0 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -53,6 +53,7 @@ set(TEST_SRC hipGraphMemsetNodeSetParams.cc hipGraphExecMemcpyNodeSetParamsFromSymbol.cc hipGraphEventRecordNodeGetEvent.cc + hipGraphEventRecordNodeSetEvent.cc hipGraphEventWaitNodeGetEvent.cc hipGraphExecMemcpyNodeSetParams.cc ) diff --git a/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc b/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc new file mode 100644 index 0000000000..9c0df129d6 --- /dev/null +++ b/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc @@ -0,0 +1,246 @@ +/* +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 : + 1) Set a different type of event using hipGraphEventRecordNodeSetEvent and + validate using hipGraphEventRecordNodeGetEvent. + 2) Add different kinds of nodes to graph and add dependencies to nodes. + Create an event record node at the end with Default flag. Set a different + type of event using hipGraphEventRecordNodeSetEvent. Instantiate and + Launch graph. Wait for the event to complete. Verify the results. + 3) Negative Scenarios + - Input node parameter is nullptr. + - Input event parameter is nullptr. + - Empty node is passed as input node. + - Memset node is passed as input node. + - Input node is an uninitialized node. + - Input event is an uninitialized event. +*/ + +#include +#include +#include + +/** + * Local Function: Set Get test + */ +static void validateEventRecordNodeSetEvent(unsigned flag) { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create events + hipEvent_t event1, event2, event_out; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreateWithFlags(&event2, flag)); + hipGraphNode_t eventrec; + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, + event1)); + // Set a different event + HIP_CHECK(hipGraphEventRecordNodeSetEvent(eventrec, event2)); + HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out)); + // validate set event and get event are same + REQUIRE(event2 == event_out); + // Free resources + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +} + +/** + * Local Function + */ +static void setEventWaitNode() { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create events + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventwait; + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, + event1)); + // Set a different event eventwait using hipGraphEventRecordNodeSetEvent + REQUIRE(hipErrorInvalidValue == + hipGraphEventRecordNodeSetEvent(eventwait, event2)); + // Free resources + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +} + +/** + * Scenario 2: Validate Change of event property in event record node. + */ +TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create events + hipEvent_t event1_start, event2_start, event1_end, event2_end; + HIP_CHECK(hipEventCreate(&event1_start)); + HIP_CHECK(hipEventCreate(&event1_end)); + HIP_CHECK(hipEventCreateWithFlags(&event2_start, hipEventDisableTiming)); + HIP_CHECK(hipEventCreateWithFlags(&event2_end, hipEventDisableTiming)); + // Create nodes + hipGraphNode_t event_start_rec, event_end_rec; + HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0, + event1_start)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0, + event1_end)); + // Create memset node + constexpr size_t Nbytes = 1024; + char *A_d; + hipGraphNode_t memset_A; + hipMemsetParams memsetParams{}; + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + // Create dependencies + // event_start_rec --> memset_A --> event_end_rec + HIP_CHECK(hipGraphAddDependencies(graph, &event_start_rec, &memset_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &event_end_rec, 1)); + + // Instantiate and launch graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + // Validate by measuring time difference between event_end_rec & + // event_start_rec + float t = 0.0f; + REQUIRE(hipSuccess == hipEventElapsedTime(&t, event1_start, + event1_end)); + REQUIRE(t > 0.0f); + // Change the event property after instantiation + HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_start_rec, event2_start)); + HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_end_rec, event2_end)); + // Launch the graph again with the new settings + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + // Validate hipEventElapsedTime() must return + // hipErrorInvalidHandle when events are created using + // hipEventDisableTiming flag. + t = 0.0f; + REQUIRE(hipErrorInvalidHandle == + hipEventElapsedTime(&t, event2_start, event2_end)); + // Free resources + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipEventDestroy(event1_start)); + HIP_CHECK(hipEventDestroy(event2_start)); + HIP_CHECK(hipEventDestroy(event1_end)); + HIP_CHECK(hipEventDestroy(event2_end)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Scenario 1: Validate Set Get test for all Event flags + */ +TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetGet") { + SECTION("Flag = hipEventDefault") { + validateEventRecordNodeSetEvent(hipEventDefault); + } + + SECTION("Flag = hipEventBlockingSync") { + validateEventRecordNodeSetEvent(hipEventBlockingSync); + } + + SECTION("Flag = hipEventDisableTiming") { + validateEventRecordNodeSetEvent(hipEventDisableTiming); + } +} + +/** + * Scenario 3: Negative Tests + */ +TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_Negative") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventrec; + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, + event1)); + SECTION("node = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeSetEvent(nullptr, + event2)); + } + + SECTION("event_out = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeSetEvent(eventrec, + nullptr)); + } + + SECTION("input node is empty node") { + hipGraphNode_t EmptyGraphNode; + HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == + hipGraphEventRecordNodeSetEvent(EmptyGraphNode, event2)); + } + + SECTION("input node is memset node") { + constexpr size_t Nbytes = 1024; + char *A_d; + hipGraphNode_t memset_A; + hipMemsetParams memsetParams{}; + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + REQUIRE(hipErrorInvalidValue == + hipGraphEventRecordNodeSetEvent(memset_A, event2)); + HIP_CHECK(hipFree(A_d)); + } + + SECTION("input node is event wait node") { + setEventWaitNode(); + } + + SECTION("input node is uninitialized node") { + hipGraphNode_t node_uninit{}; + REQUIRE(hipErrorInvalidValue == + hipGraphEventRecordNodeSetEvent(node_uninit, event2)); + } + + SECTION("input event is uninitialized") { + hipEvent_t event_uninit{}; + REQUIRE(hipErrorInvalidValue == + hipGraphEventRecordNodeSetEvent(eventrec, event_uninit)); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +}