diff --git a/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc index 4c544a3f68..a6a028e604 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -29,12 +29,32 @@ Functional :: 1) Create Node and destroy the node 2) Create graph with dependencies and destroy one of the dependency node before executing the graph. +3) Create a graph with N nodes and (N-1) dependencies between them as shown + below. Start destroying the nodes in iteration from left. In each iteration + verify the number of nodes and dependencies using hipGraphGetNodes and + hipGraphGetEdges. + Node1-->Node2-->Node3->...................->NodeN +4) Create a graph with N nodes and (N-1) dependencies between them as shown + above. Clone the graph. Start destroying the nodes in iteration from left + in the cloned graph. In each iteration verify the number of nodes and + dependencies using hipGraphGetNodes and hipGraphGetEdges. Once all nodes + in the cloned graph are deleted, verify the number of nodes in the original + graph are intact. +5) Create a graph1 with N nodes and (N-1) dependencies between them as shown + above. Create another empty graph0. Add graph1 as child node to graph0. + Delete the child node in graph0. Verify that the nodes in graph1 are still + intact after deleting the child node using hipGraphGetNodes and hipGraphGetEdges. */ #include #include #include +#define NUM_OF_DUMMY_NODES 8 + +static __global__ void dummyKernel() { + return; +} /* This test covers the negative scenarios of hipGraphDestroyNode API */ @@ -137,3 +157,142 @@ TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { HIP_CHECK(hipStreamDestroy(streamForGraph)); HIP_CHECK(hipGraphDestroy(graph)); } + +/** + * Functional Test to test hipGraphDestroyNode using hipGraphGetNodes + * and hipGraphGetEdges APIs. + */ +TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep") { + hipGraph_t graph; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + // Create dependencies between nodes + for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) { + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i-1], + &kernelnode[i], 1)); + } + // Start destroying nodes from 0 + size_t numOfNodes = 0, numOfDep = 0; + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + // destroy node i + HIP_CHECK(hipGraphDestroyNode(kernelnode[i])); + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes)); + REQUIRE(numOfNodes == (NUM_OF_DUMMY_NODES - i - 1)); + HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numOfDep)); + REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2)); + } + HIP_CHECK(hipGraphDestroyNode(kernelnode[NUM_OF_DUMMY_NODES-1])); + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes)); + REQUIRE(numOfNodes == 0); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test to test hipGraphDestroyNode using hipGraphGetNodes + * and hipGraphGetEdges APIs on a cloned graph + */ +TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph") { + hipGraph_t graph, clonedgraph; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&clonedgraph, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + // Create dependencies between nodes + for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) { + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i-1], + &kernelnode[i], 1)); + } + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Start destroying nodes from 0 and validate number of nodes in + // cloned graph + size_t numOfNodes = 0, numOfDep = 0; + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + hipGraphNode_t node; + // destroy node i + HIP_CHECK(hipGraphNodeFindInClone(&node, kernelnode[i], clonedgraph)); + HIP_CHECK(hipGraphDestroyNode(node)); + HIP_CHECK(hipGraphGetNodes(clonedgraph, nullptr, &numOfNodes)); + REQUIRE(numOfNodes == (NUM_OF_DUMMY_NODES - i - 1)); + HIP_CHECK(hipGraphGetEdges(clonedgraph, nullptr, nullptr, &numOfDep)); + REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2)); + } + // Verify the number of nodes in original graph + numOfNodes = 0; + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numOfNodes)); + REQUIRE(numOfNodes == NUM_OF_DUMMY_NODES); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test to test hipGraphDestroyNode on child node using + * hipGraphGetNodes and hipGraphGetEdges APIs on a cloned graph. + */ +TEST_CASE("Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode") { + hipGraph_t graph0, graph1; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES], childGraphNode; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph0, 0)); + HIP_CHECK(hipGraphCreate(&graph1, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph0, nullptr, + 0, &kernelNodeParams[i])); + } + // Create dependencies between nodes + for (int i = 1; i < NUM_OF_DUMMY_NODES; i++) { + HIP_CHECK(hipGraphAddDependencies(graph0, &kernelnode[i-1], + &kernelnode[i], 1)); + } + // Create child node and add it to graph1 + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph1, + nullptr, 0, graph0)); + // delete the child node from graph1 + HIP_CHECK(hipGraphDestroyNode(childGraphNode)); + // Start destroying nodes from 0 + size_t numOfNodes = 0, numOfDep = 0; + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + // destroy node i + HIP_CHECK(hipGraphDestroyNode(kernelnode[i])); + HIP_CHECK(hipGraphGetNodes(graph0, nullptr, &numOfNodes)); + REQUIRE(numOfNodes == (NUM_OF_DUMMY_NODES - i - 1)); + HIP_CHECK(hipGraphGetEdges(graph0, nullptr, nullptr, &numOfDep)); + REQUIRE(numOfDep == (NUM_OF_DUMMY_NODES - i - 2)); + } + HIP_CHECK(hipGraphGetNodes(graph1, nullptr, &numOfNodes)); + REQUIRE(numOfNodes == 0); + HIP_CHECK(hipGraphDestroy(graph0)); + HIP_CHECK(hipGraphDestroy(graph1)); +}