SWDEV-306122 - [catch2][dtest] hipGraph tests for GetNodes/GetRootNodes (#2497)
Change-Id: I5e9fafd0802585b6b33c391d7281151010ffc1c3
[ROCm/hip-tests commit: 3a4fde33e2]
This commit is contained in:
committed by
GitHub
parent
26d1b29634
commit
1ca2bb099a
@@ -1,5 +1,5 @@
|
|||||||
/*
|
/*
|
||||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||||
of this software and associated documentation files (the "Software"), to deal
|
of this software and associated documentation files (the "Software"), to deal
|
||||||
in the Software without restriction, including without limitation the rights
|
in the Software without restriction, including without limitation the rights
|
||||||
@@ -18,13 +18,20 @@ THE SOFTWARE.
|
|||||||
*/
|
*/
|
||||||
|
|
||||||
/**
|
/**
|
||||||
Testcase Scenarios :
|
Testcase Scenarios
|
||||||
|
------------------
|
||||||
Functional ::
|
Functional ::
|
||||||
1) Add nodes to graph and get nodes. Verify the added nodes are present in returned list.
|
1) Add nodes to graph and get nodes. Verify the added nodes are present in returned list.
|
||||||
2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph.
|
2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph.
|
||||||
3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes
|
3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes
|
||||||
will be set to NULL, and the number of nodes actually obtained will be returned in numNodes.
|
will be set to NULL, and the number of nodes actually obtained will be returned in numNodes.
|
||||||
|
|
||||||
|
Argument Validation ::
|
||||||
|
1) Pass graph as nullptr and verify api returns error code.
|
||||||
|
2) Pass numNodes as nullptr and other params as valid values. Expect api to return error code.
|
||||||
|
3) When there are no nodes in graph, expect numNodes to be set to zero.
|
||||||
|
4) Pass numNodes less than actual number of nodes. Expect api to populate requested number of node entries
|
||||||
|
and does update numNodes.
|
||||||
*/
|
*/
|
||||||
|
|
||||||
#include <hip_test_common.hh>
|
#include <hip_test_common.hh>
|
||||||
@@ -32,7 +39,7 @@ will be set to NULL, and the number of nodes actually obtained will be returned
|
|||||||
#include <hip_test_kernels.hh>
|
#include <hip_test_kernels.hh>
|
||||||
|
|
||||||
/**
|
/**
|
||||||
* Functional Test for API fetching node list
|
* Functional Test for hipGraphGetNodes API fetching node list
|
||||||
*/
|
*/
|
||||||
TEST_CASE("Unit_hipGraphGetNodes_Functional") {
|
TEST_CASE("Unit_hipGraphGetNodes_Functional") {
|
||||||
constexpr size_t N = 1024;
|
constexpr size_t N = 1024;
|
||||||
@@ -131,3 +138,82 @@ TEST_CASE("Unit_hipGraphGetNodes_Functional") {
|
|||||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||||
free(nodes);
|
free(nodes);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
/**
|
||||||
|
* Test performs api parameter validation by passing various values
|
||||||
|
* as input and output parameters and validates the behavior.
|
||||||
|
* Test will include both negative and positive scenarios.
|
||||||
|
*/
|
||||||
|
TEST_CASE("Unit_hipGraphGetNodes_ParamValidation") {
|
||||||
|
hipStream_t stream{nullptr};
|
||||||
|
hipGraph_t graph{nullptr};
|
||||||
|
constexpr unsigned blocks = 512;
|
||||||
|
constexpr unsigned threadsPerBlock = 256;
|
||||||
|
constexpr size_t N = 1000000;
|
||||||
|
size_t Nbytes = N * sizeof(float), numNodes{};
|
||||||
|
float *A_d, *C_d;
|
||||||
|
float *A_h, *C_h;
|
||||||
|
A_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||||
|
C_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||||
|
REQUIRE(A_h != nullptr);
|
||||||
|
REQUIRE(C_h != nullptr);
|
||||||
|
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||||
|
HIP_CHECK(hipMalloc(&C_d, Nbytes));
|
||||||
|
REQUIRE(A_d != nullptr);
|
||||||
|
REQUIRE(C_d != nullptr);
|
||||||
|
|
||||||
|
HIP_CHECK(hipStreamCreate(&stream));
|
||||||
|
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
|
||||||
|
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||||
|
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
|
||||||
|
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||||
|
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
|
||||||
|
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||||
|
HIP_CHECK(hipStreamEndCapture(stream, &graph));
|
||||||
|
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
|
||||||
|
INFO("Num of nodes returned by GetNodes : " << numNodes);
|
||||||
|
|
||||||
|
int numBytes = sizeof(hipGraphNode_t) * numNodes;
|
||||||
|
hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
|
||||||
|
REQUIRE(nodes != nullptr);
|
||||||
|
|
||||||
|
SECTION("graph as nullptr") {
|
||||||
|
hipError_t ret = hipGraphGetNodes(nullptr, nodes, &numNodes);
|
||||||
|
REQUIRE(ret == hipErrorInvalidValue);
|
||||||
|
}
|
||||||
|
|
||||||
|
SECTION("numNodes as nullptr") {
|
||||||
|
hipError_t ret = hipGraphGetNodes(graph, nodes, nullptr);
|
||||||
|
REQUIRE(ret == hipErrorInvalidValue);
|
||||||
|
}
|
||||||
|
|
||||||
|
SECTION("no nodes in graph") {
|
||||||
|
hipGraph_t emptyGraph{};
|
||||||
|
HIP_CHECK(hipGraphCreate(&emptyGraph, 0));
|
||||||
|
HIP_CHECK(hipGraphGetNodes(emptyGraph, nullptr, &numNodes));
|
||||||
|
REQUIRE(numNodes == 0);
|
||||||
|
}
|
||||||
|
|
||||||
|
SECTION("numNodes less than actual number of nodes") {
|
||||||
|
size_t numPartNodes = numNodes - 1;
|
||||||
|
hipGraphNodeType nodeType;
|
||||||
|
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numPartNodes));
|
||||||
|
|
||||||
|
// verify numPartNodes is unchanged
|
||||||
|
REQUIRE(numPartNodes == numNodes - 1);
|
||||||
|
// verify partial node list returned has valid nodes
|
||||||
|
for (size_t i = 0; i < numPartNodes; i++) {
|
||||||
|
HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType));
|
||||||
|
REQUIRE(nodeType >= 0);
|
||||||
|
REQUIRE(nodeType < hipGraphNodeTypeCount);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
HIP_CHECK(hipGraphDestroy(graph));
|
||||||
|
HIP_CHECK(hipStreamDestroy(stream));
|
||||||
|
free(A_h);
|
||||||
|
free(C_h);
|
||||||
|
free(nodes);
|
||||||
|
HIP_CHECK(hipFree(A_d));
|
||||||
|
HIP_CHECK(hipFree(C_d));
|
||||||
|
}
|
||||||
|
|||||||
@@ -1,5 +1,5 @@
|
|||||||
/*
|
/*
|
||||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||||
of this software and associated documentation files (the "Software"), to deal
|
of this software and associated documentation files (the "Software"), to deal
|
||||||
in the Software without restriction, including without limitation the rights
|
in the Software without restriction, including without limitation the rights
|
||||||
@@ -18,13 +18,21 @@ THE SOFTWARE.
|
|||||||
*/
|
*/
|
||||||
|
|
||||||
/**
|
/**
|
||||||
Testcase Scenarios :
|
Testcase Scenarios
|
||||||
|
------------------
|
||||||
Functional ::
|
Functional ::
|
||||||
1) Add nodes to graph with and without dependencies, verify the api returns list of
|
1) Add nodes to graph with and without dependencies, verify the api returns list of
|
||||||
root nodes (i.e., nodes without dependencies).
|
root nodes (i.e., nodes without dependencies).
|
||||||
2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph.
|
2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph.
|
||||||
3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in
|
3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in
|
||||||
nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes.
|
nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes.
|
||||||
|
|
||||||
|
Argument Validation ::
|
||||||
|
1) Pass graph as nullptr and verify api returns error code.
|
||||||
|
2) Pass numRootNodes as nullptr and other params as valid values. Expect api to return error code.
|
||||||
|
3) When there are no nodes in graph, expect numRootNodes to be set to zero.
|
||||||
|
4) Pass numRootNodes less than actual number of nodes. Expect api to populate requested number of node entries
|
||||||
|
and does update numRootNodes.
|
||||||
*/
|
*/
|
||||||
|
|
||||||
#include <hip_test_common.hh>
|
#include <hip_test_common.hh>
|
||||||
@@ -41,6 +49,7 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") {
|
|||||||
constexpr auto threadsPerBlock = 256;
|
constexpr auto threadsPerBlock = 256;
|
||||||
constexpr auto addlEntries = 5;
|
constexpr auto addlEntries = 5;
|
||||||
hipGraph_t graph;
|
hipGraph_t graph;
|
||||||
|
|
||||||
hipGraphNode_t memcpyNode, kernelNode;
|
hipGraphNode_t memcpyNode, kernelNode;
|
||||||
hipKernelNodeParams kernelNodeParams{};
|
hipKernelNodeParams kernelNodeParams{};
|
||||||
hipStream_t streamForGraph;
|
hipStream_t streamForGraph;
|
||||||
@@ -131,3 +140,100 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") {
|
|||||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||||
free(rootnodes);
|
free(rootnodes);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
/**
|
||||||
|
* Test performs api parameter validation by passing various values
|
||||||
|
* as input and output parameters and validates the behavior.
|
||||||
|
* Test will include both negative and positive scenarios.
|
||||||
|
*/
|
||||||
|
TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") {
|
||||||
|
hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr};
|
||||||
|
hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent;
|
||||||
|
hipGraph_t graph{nullptr};
|
||||||
|
constexpr unsigned blocks = 512;
|
||||||
|
constexpr unsigned threadsPerBlock = 256;
|
||||||
|
constexpr size_t N = 1000000;
|
||||||
|
size_t Nbytes = N * sizeof(float), numRootNodes{};
|
||||||
|
float *A_d, *C_d;
|
||||||
|
float *A_h, *C_h;
|
||||||
|
A_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||||
|
C_h = reinterpret_cast<float*>(malloc(Nbytes));
|
||||||
|
REQUIRE(A_h != nullptr);
|
||||||
|
REQUIRE(C_h != nullptr);
|
||||||
|
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||||
|
HIP_CHECK(hipMalloc(&C_d, Nbytes));
|
||||||
|
REQUIRE(A_d != nullptr);
|
||||||
|
REQUIRE(C_d != nullptr);
|
||||||
|
|
||||||
|
HIP_CHECK(hipStreamCreate(&stream1));
|
||||||
|
HIP_CHECK(hipStreamCreate(&stream2));
|
||||||
|
HIP_CHECK(hipStreamCreate(&mstream));
|
||||||
|
HIP_CHECK(hipEventCreate(&memsetEvent1));
|
||||||
|
HIP_CHECK(hipEventCreate(&memsetEvent2));
|
||||||
|
HIP_CHECK(hipEventCreate(&forkStreamEvent));
|
||||||
|
HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal));
|
||||||
|
HIP_CHECK(hipEventRecord(forkStreamEvent, mstream));
|
||||||
|
HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0));
|
||||||
|
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
|
||||||
|
HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1));
|
||||||
|
HIP_CHECK(hipEventRecord(memsetEvent1, stream1));
|
||||||
|
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2));
|
||||||
|
HIP_CHECK(hipEventRecord(memsetEvent2, stream2));
|
||||||
|
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0));
|
||||||
|
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0));
|
||||||
|
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream));
|
||||||
|
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
|
||||||
|
dim3(threadsPerBlock), 0, mstream, A_d, C_d, N);
|
||||||
|
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream));
|
||||||
|
HIP_CHECK(hipStreamEndCapture(mstream, &graph));
|
||||||
|
HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes));
|
||||||
|
INFO("Num of nodes returned by GetRootNodes : " << numRootNodes);
|
||||||
|
int numBytes = sizeof(hipGraphNode_t) * numRootNodes;
|
||||||
|
hipGraphNode_t* nodes = reinterpret_cast<hipGraphNode_t *>(malloc(numBytes));
|
||||||
|
REQUIRE(nodes != nullptr);
|
||||||
|
|
||||||
|
SECTION("graph as nullptr") {
|
||||||
|
hipError_t ret = hipGraphGetRootNodes(nullptr, nodes, &numRootNodes);
|
||||||
|
REQUIRE(ret == hipErrorInvalidValue);
|
||||||
|
}
|
||||||
|
|
||||||
|
SECTION("numRootNodes as nullptr") {
|
||||||
|
hipError_t ret = hipGraphGetRootNodes(graph, nodes, nullptr);
|
||||||
|
REQUIRE(ret == hipErrorInvalidValue);
|
||||||
|
}
|
||||||
|
|
||||||
|
SECTION("no nodes in graph") {
|
||||||
|
hipGraph_t emptyGraph{};
|
||||||
|
HIP_CHECK(hipGraphCreate(&emptyGraph, 0));
|
||||||
|
HIP_CHECK(hipGraphGetRootNodes(emptyGraph, nullptr, &numRootNodes));
|
||||||
|
REQUIRE(numRootNodes == 0);
|
||||||
|
}
|
||||||
|
|
||||||
|
SECTION("numRootNodes less than actual number of nodes") {
|
||||||
|
size_t numPartNodes = numRootNodes - 1;
|
||||||
|
hipGraphNodeType nodeType;
|
||||||
|
HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numPartNodes));
|
||||||
|
|
||||||
|
// verify numPartNodes is unchanged
|
||||||
|
REQUIRE(numPartNodes == numRootNodes - 1);
|
||||||
|
// verify partial node list returned has valid nodes
|
||||||
|
for (size_t i = 0; i < numPartNodes; i++) {
|
||||||
|
HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType));
|
||||||
|
REQUIRE(nodeType >= 0);
|
||||||
|
REQUIRE(nodeType < hipGraphNodeTypeCount);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
HIP_CHECK(hipGraphDestroy(graph));
|
||||||
|
HIP_CHECK(hipStreamDestroy(mstream));
|
||||||
|
HIP_CHECK(hipStreamDestroy(stream1));
|
||||||
|
HIP_CHECK(hipStreamDestroy(stream2));
|
||||||
|
HIP_CHECK(hipEventDestroy(forkStreamEvent));
|
||||||
|
HIP_CHECK(hipEventDestroy(memsetEvent1));
|
||||||
|
HIP_CHECK(hipEventDestroy(memsetEvent2));
|
||||||
|
free(A_h);
|
||||||
|
free(C_h);
|
||||||
|
free(nodes);
|
||||||
|
HIP_CHECK(hipFree(A_d));
|
||||||
|
HIP_CHECK(hipFree(C_d));
|
||||||
|
}
|
||||||
|
|||||||
Reference in New Issue
Block a user