From 1e15bfca2317bc7893c9e4a134ced61485bbe3ef Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Fri, 29 Dec 2023 14:29:03 +0000 Subject: [PATCH] EXSWHTEC-383 - Implement tests for hipGraphAddNode API #456 Change-Id: I91762d23c4a05acc68463402e76cee4e26d27b57 --- catch/unit/graph/CMakeLists.txt | 1 + catch/unit/graph/hipGraphAddNode.cc | 553 ++++++++++++++++++++++++++++ 2 files changed, 554 insertions(+) create mode 100644 catch/unit/graph/hipGraphAddNode.cc diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index d6519375a4..6a077cc80e 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -157,6 +157,7 @@ set(TEST_SRC hipDeviceSetGraphMemAttribute.cc hipDeviceGetGraphMemAttribute.cc hipDeviceGraphMemTrim.cc + hipGraphAddNode.cc ) add_custom_target(add_Kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/add_Kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../graph/add_Kernel.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) diff --git a/catch/unit/graph/hipGraphAddNode.cc b/catch/unit/graph/hipGraphAddNode.cc new file mode 100644 index 0000000000..645fd0716f --- /dev/null +++ b/catch/unit/graph/hipGraphAddNode.cc @@ -0,0 +1,553 @@ +/* +Copyright (c) 2023 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 +#include +#include +#include +#include +#include + +#include "graph_memset_node_test_common.hh" +#include "graph_tests_common.hh" + +#pragma clang diagnostic ignored "-Wunused-parameter" + +/** + * @addtogroup hipGraphAddNode hipGraphAddNode + * @{ + * @ingroup GraphTest + * `hipGraphAddNode(hipGraphNode_t *pGraphNode, hipGraph_t graph, const hipGraphNode_t + * *pDependencies, size_t numDependencies, hipGraphNodeParams *nodeParams)` - Creates a node and + * adds it to a graph + */ + +static constexpr size_t N = 1024; + +static void callbackfunc(void* A_h) { + int* A = reinterpret_cast(A_h); + for (int i = 0; i < N; i++) { + A[i] = i; + } +} + +static void __global__ vector_square(int* A_d) { + for (int i = 0; i < N; i++) { + A_d[i] = A_d[i] * A_d[i]; + } +} + +/** + * Test Description + * ------------------------ + * - Verify that all elements of destination memory are set to the correct value. + * The test is repeated for all valid element sizes(1, 2, 4), and several allocations of different + * height and width, both on host and device. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEMPLATE_TEST_CASE("Unit_hipGraphAddNodeTypeMemset_Positive_Basic", "", uint8_t, uint16_t, + uint32_t) { + const auto f = [](hipMemsetParams* params) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t node = nullptr; + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeMemset; + node_params.memset.dst = params->dst; + node_params.memset.elementSize = params->elementSize; + node_params.memset.width = params->width; + node_params.memset.height = params->height; + node_params.memset.pitch = params->pitch; + node_params.memset.value = params->value; + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; + }; + + GraphMemsetNodeCommonPositive(f); +} + +/** + * Test Description + * ------------------------ + * - Verify that kernel node added with hipGraphAddNode executes correctly and does the square of + * values in the device array. The result is copied to host and verified. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeKernel_Positive_Basic") { + constexpr size_t allocation_size = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + + int* A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, &A_h, &B_h, nullptr, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t memcpyH2D_A, memcpyD2H_B; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, allocation_size, + hipMemcpyHostToDevice)); + + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeKernel; + void* kernel_args[] = {&A_d}; + node_params.kernel.func = reinterpret_cast(vector_square); + node_params.kernel.gridDim = dim3(1); + node_params.kernel.blockDim = dim3(1); + node_params.kernel.sharedMemBytes = 0; + node_params.kernel.kernelParams = reinterpret_cast(kernel_args); + node_params.kernel.extra = nullptr; + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, nullptr, 0, B_h, A_d, allocation_size, + hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &node, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &node, &memcpyD2H_B, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (B_h[i] != (A_h[i] * A_h[i])) { + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Test Description + * ------------------------ + * - Verify that host node added with hipGraphAddNode executes correctly and sets values of host + * array. The result is verified. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeHost_Positive_Basic") { + constexpr size_t allocation_size = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int* A_h = (int*)malloc(allocation_size); + std::fill_n(A_h, N, 0); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeHost; + node_params.host.fn = callbackfunc; + node_params.host.userData = A_h; + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (A_h[i] != static_cast(i)) { + REQUIRE(false); + } + } + + free(A_h); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Test Description + * ------------------------ + * - Verify that when graph is created and childgraph node is added with hipGraphAddNode, the + * childgraph executes correctly. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeChildGraph_Positive_Basic") { + constexpr size_t allocation_size = N * sizeof(int); + hipGraph_t graph, childgraph; + hipGraphExec_t graphExec; + + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + for (size_t i = 0; i < N; i++) { + B_h[i] = i; + } + + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&childgraph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, childgraph, nullptr, 0, B_d, B_h, allocation_size, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph, nullptr, 0, A_h, B_d, allocation_size, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, 0, C_d, C_h, allocation_size, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, 0, A_h, C_d, allocation_size, + hipMemcpyDeviceToHost)); + + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeGraph; + node_params.graph.graph = childgraph; + HIP_CHECK(hipGraphAddNode(&childGraphNode1, graph, nullptr, 0, &node_params)); + + HIP_CHECK(hipGraphAddDependencies(childgraph, &memcpyH2D_B, &memcpyH2D_A, 1)); + + // Instantiate and launch the childgraph + HIP_CHECK(hipGraphInstantiate(&graphExec, childgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (B_h[i] != A_h[i]) { + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(childgraph)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + + +static hipError_t MemcpyType3DWrapper(PtrVariant dst_ptr, hipPos dst_pos, PtrVariant src_ptr, + hipPos src_pos, hipExtent extent, hipMemcpyKind kind, + hipStream_t stream = nullptr) { + auto parms = GetMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind); + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeMemcpy; + memset(&node_params.memcpy, 0, sizeof(hipMemcpyNodeParams)); + node_params.memcpy.copyParams = parms; + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; +} + +/** + * Test Description + * ------------------------ + * - Verify basic API behavior. A Memcpy node is created using hipGraphAddNode with parameters + * set according to the test run, after which the graph is run and the memcpy results are verified. + * The test is run for all possible memcpy directions, with both the corresponding memcpy + * kind and hipMemcpyDefault, as well as half page and full page allocation sizes. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeMemcpy_Positive_Basic") { + constexpr bool async = false; + + SECTION("Device to host") { Memcpy3DDeviceToHostShell(MemcpyType3DWrapper); } + + SECTION("Device to host with default kind") { + Memcpy3DDeviceToHostShell(MemcpyType3DWrapper); + } + + SECTION("Host to device") { Memcpy3DHostToDeviceShell(MemcpyType3DWrapper); } + + SECTION("Host to device with default kind") { + Memcpy3DHostToDeviceShell(MemcpyType3DWrapper); + } + + SECTION("Host to host") { Memcpy3DHostToHostShell(MemcpyType3DWrapper); } + + SECTION("Host to host with default kind") { Memcpy3DHostToHostShell(MemcpyType3DWrapper); } + + SECTION("Device to device") { + SECTION("Peer access enabled") { + Memcpy3DDeviceToDeviceShell(MemcpyType3DWrapper); + } + SECTION("Peer access disabled") { + Memcpy3DDeviceToDeviceShell(MemcpyType3DWrapper); + } + } + + SECTION("Device to device with default kind") { + SECTION("Peer access enabled") { + Memcpy3DDeviceToDeviceShell(MemcpyType3DWrapper); + } + SECTION("Peer access disabled") { + Memcpy3DDeviceToDeviceShell(MemcpyType3DWrapper); + } + } + + SECTION("Array from/to Host") { Memcpy3DArrayHostShell(MemcpyType3DWrapper); } + +#if HT_NVIDIA // Disabled on AMD due to defect - EXSWHTEC-220 + SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell(MemcpyType3DWrapper); } +#endif +} + + +/** + * Test Description + * ------------------------ + * - Verify basic API functionality where one event record node is added to graph with + * hipGraphAddNode and its correct behavior is verified. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeEventRecord_Positive_Basic") { + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeEventRecord; + node_params.eventRecord.event = event; + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + // Wait for event + HIP_CHECK(hipEventSynchronize(event)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Test Description + * ------------------------ + * - Verify basic API functionality where one event record and one event wait nodes are added to + * graph with hipGraphAddNode and their correct behavior is verified. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeEventWait_Positive_Basic") { + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + hipGraphNode_t event_rec_node, event_wait_node; + + // Create a event record node in graph + hipGraphNodeParams rec_node_params = {}; + rec_node_params.type = hipGraphNodeTypeEventRecord; + rec_node_params.eventRecord.event = event; + HIP_CHECK(hipGraphAddNode(&event_rec_node, graph, nullptr, 0, &rec_node_params)); + + // Create a event wait node in graph + hipGraphNodeParams wait_node_params = {}; + rec_node_params.type = hipGraphNodeTypeWaitEvent; + rec_node_params.eventWait.event = event; + HIP_CHECK(hipGraphAddNode(&event_wait_node, graph, &event_rec_node, 1, &wait_node_params)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Test Description + * ------------------------ + * - Test to verify basic API functionality when memalloc and memfree nodes are added with + * hipGraphAddNode. Verify that memory is allocated correctly and graph behaves as expected when + * free node is added to the same graph. + * Test source + * ------------------------ + * - /unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNodeTypeMemAlloc_Positive_Basic") { + constexpr size_t allocation_size = N * sizeof(int); + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t alloc_node; + hipGraphNodeParams alloc_node_params = {}; + alloc_node_params.type = hipGraphNodeTypeMemAlloc; + memset(&alloc_node_params.alloc, 0, sizeof(hipMemAllocNodeParams)); + alloc_node_params.alloc.bytesize = allocation_size; + alloc_node_params.alloc.poolProps.allocType = hipMemAllocationTypePinned; + alloc_node_params.alloc.poolProps.location.id = 0; + alloc_node_params.alloc.poolProps.location.type = hipMemLocationTypeDevice; + HIP_CHECK(hipGraphAddNode(&alloc_node, graph, nullptr, 0, &alloc_node_params)); + + REQUIRE(alloc_node_params.alloc.dptr != nullptr); + int* A_d = reinterpret_cast(alloc_node_params.alloc.dptr); + + hipGraphNode_t free_node; + hipGraphNodeParams free_node_params = {}; + free_node_params.type = hipGraphNodeTypeMemFree; + free_node_params.free.dptr = A_d; + HIP_CHECK(hipGraphAddNode(&free_node, graph, &alloc_node, 1, &free_node_params)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipGraphAddNode behavior with invalid arguments: + * -# Nullptr graph + * -# Nullptr graph node + * -# Invalid numDependencies for null list of dependencies + * -# Node in dependency is from different graph + * -# Invalid numNodes + * -# Duplicate node in dependencies + * -# Nullptr params + * -# params type is invalid + * Test source + * ------------------------ + * - /unit/graph/hipGraphAddNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddNode_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + node_params.type = hipGraphNodeTypeEventRecord; + node_params.eventRecord.event = event; + + GraphAddNodeCommonNegativeTests(std::bind(hipGraphAddNode, _1, _2, _3, _4, &node_params), graph); + + SECTION("params == nullptr") { + HIP_CHECK_ERROR(hipGraphAddNode(&node, graph, nullptr, 0, nullptr), hipErrorInvalidValue); + } + + SECTION("params type is invalid") { + node_params.type = static_cast(0x20); + HIP_CHECK_ERROR(hipGraphAddNode(&node, graph, nullptr, 0, &node_params), hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); +}