AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-366390 - [catch2][dtest] Adding additional functional tests for hipGraphDestroyNode() API (#101)

Browse Source 
Change-Id: Ia906f9a1c409c1d694be9089f75e939729e6103b


[ROCm/hip-tests commit: 70c61d55a4]
This commit is contained in:
ROCm CI Service Account
2023-01-07 01:30:07 +05:30
committed by GitHub
parent 1af3cd3067
commit 9bc91a13b2
1 changed files with 160 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/graph/hipGraphDestroyNode.cc
+160 -1
View File
@@ -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 <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#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<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(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<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(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<void *>(dummyKernel);
kernelNodeParams[i].gridDim = dim3(1);
kernelNodeParams[i].blockDim = dim3(1);
kernelNodeParams[i].sharedMemBytes = 0;
kernelNodeParams[i].kernelParams = reinterpret_cast<void**>(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));
}