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

EXSWHTEC-178 - Implement tests for Graph Event Node APIs #44

Browse Source 
Change-Id: Id3f569d94d347af2f5e27513fa01c5a1e8e30fd9


[ROCm/hip-tests commit: 575e4cc93e]
This commit is contained in:
Nives Vukovic
2023-11-24 21:42:07 +05:30
committed by Maneesh Gupta
parent 639ba1e5f3
commit bb352913bb
8 changed files with 472 additions and 441 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/graph/hipGraphAddEventWaitNode.cc
+99 -81
View File
@@ -32,20 +32,25 @@ both graphs.
(100 times).
4) Execute scenario 2 with stream1 = stream2.
5) Repeat scenario 2 for different event flags.
6) Negative Scenarios
6) Validate that no error is reported when numDeps <= dependencies length
7) Negative Scenarios
- Pass input node parameter as nullptr.
- Pass input graph parameter as nullptr.
- Pass input dependency parameter as nullptr.
- Node in dependency is from different graph
- Invalid numNodes
- Duplicate node in dependencies
- Pass input event parameter as nullptr.
- Pass uninitialized input graph parameter.
- Pass uninitialized input event parameter.
*/
#include <functional>
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#define LEN 512
#include "graph_tests_common.hh"
/**
* Scenario 1
@@ -60,13 +65,10 @@ TEST_CASE("Unit_hipGraphAddEventWaitNode_Functional_Simple") {
HIP_CHECK(hipEventCreate(&event));
hipGraphNode_t event_rec_node, event_wait_node;
// Create a event record node in graph
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph, nullptr, 0,
event));
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph, nullptr, 0, event));
// Create a event wait node in graph
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0,
event));
HIP_CHECK(hipGraphAddDependencies(graph, &event_rec_node,
&event_wait_node, 1));
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0, event));
HIP_CHECK(hipGraphAddDependencies(graph, &event_rec_node, &event_wait_node, 1));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
@@ -80,13 +82,14 @@ TEST_CASE("Unit_hipGraphAddEventWaitNode_Functional_Simple") {
/**
* Local Function
*/
static void validate_hipGraphAddEventWaitNode_internodedep(int test,
int nstep, unsigned flag = hipEventDefault) {
size_t memsize = LEN * sizeof(int);
static void validate_hipGraphAddEventWaitNode_internodedep(int test, int nstep,
unsigned flag = hipEventDefault) {
constexpr size_t N = 1024;
size_t memsize = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, LEN);
size_t NElem{LEN};
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
size_t NElem{N};
hipGraph_t graph1, graph2;
hipStream_t streamForGraph1, streamForGraph2;
hipGraphExec_t graphExec1, graphExec2;
@@ -114,68 +117,57 @@ static void validate_hipGraphAddEventWaitNode_internodedep(int test,
HIP_CHECK(hipMalloc(&out_d_g1, memsize));
HIP_CHECK(hipMalloc(&out_d_g2, memsize));
// Initialize host buffer
for (uint32_t i = 0; i < LEN; i++) {
for (uint32_t i = 0; i < N; i++) {
inp_h[i] = i;
out_h_g1[i] = 0;
out_h_g2[i] = 0;
}
// Graph1 creation ...........
// Create event1 record node in graph1
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph1, nullptr, 0,
event1));
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph1, nullptr, 0, event1));
// Create memcpy and kernel nodes for graph1
hipGraphNode_t memcpyH2D, memcpyD2H_1, kernelnode_1;
hipKernelNodeParams kernelNodeParams1{};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d,
inp_h, memsize, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph1, nullptr, 0,
out_h_g1, out_d_g1, memsize, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d, inp_h, memsize,
hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph1, nullptr, 0, out_h_g1, out_d_g1, memsize,
hipMemcpyDeviceToHost));
void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast<void *>(&NElem)};
kernelNodeParams1.func =
reinterpret_cast<void *>(HipTest::vector_square<int>);
void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast<void*>(&NElem)};
kernelNodeParams1.func = reinterpret_cast<void*>(HipTest::vector_square<int>);
kernelNodeParams1.gridDim = dim3(blocks);
kernelNodeParams1.blockDim = dim3(threadsPerBlock);
kernelNodeParams1.sharedMemBytes = 0;
kernelNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph1, nullptr, 0,
&kernelNodeParams1));
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph1, nullptr, 0, &kernelNodeParams1));
// Create dependencies for graph1
HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D,
&event_rec_node, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &event_rec_node,
&kernelnode_1, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode_1,
&memcpyD2H_1, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D, &event_rec_node, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &event_rec_node, &kernelnode_1, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode_1, &memcpyD2H_1, 1));
// Graph2 creation ...........
// Create event1 record node in graph2
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0,
event1));
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0, event1));
// Create memcpy and kernel nodes for graph2
hipGraphNode_t memcpyD2H_2, kernelnode_2;
hipKernelNodeParams kernelNodeParams2{};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph2, nullptr, 0,
out_h_g2, out_d_g2, memsize, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph2, nullptr, 0, out_h_g2, out_d_g2, memsize,
hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast<void *>(&NElem)};
kernelNodeParams2.func =
reinterpret_cast<void *>(HipTest::vector_cubic<int>);
void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast<void*>(&NElem)};
kernelNodeParams2.func = reinterpret_cast<void*>(HipTest::vector_cubic<int>);
kernelNodeParams2.gridDim = dim3(blocks);
kernelNodeParams2.blockDim = dim3(threadsPerBlock);
kernelNodeParams2.sharedMemBytes = 0;
kernelNodeParams2.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams2.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph2, nullptr, 0,
&kernelNodeParams2));
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph2, nullptr, 0, &kernelNodeParams2));
// Create dependencies for graph2
HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node,
&kernelnode_2, 1));
HIP_CHECK(hipGraphAddDependencies(graph2, &kernelnode_2,
&memcpyD2H_2, 1));
HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node, &kernelnode_2, 1));
HIP_CHECK(hipGraphAddDependencies(graph2, &kernelnode_2, &memcpyD2H_2, 1));
// Instantiate and launch the graphs
HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
@@ -187,16 +179,16 @@ static void validate_hipGraphAddEventWaitNode_internodedep(int test,
HIP_CHECK(hipStreamSynchronize(streamForGraph2));
// Validate output
bool btestPassed1 = true;
for (uint32_t i = 0; i < LEN; i++) {
if (out_h_g1[i] != (inp_h[i]*inp_h[i])) {
for (uint32_t i = 0; i < N; i++) {
if (out_h_g1[i] != (inp_h[i] * inp_h[i])) {
btestPassed1 = false;
break;
}
}
REQUIRE(btestPassed1 == true);
bool btestPassed2 = true;
for (uint32_t i = 0; i < LEN; i++) {
if (out_h_g2[i] != (inp_h[i]*inp_h[i]*inp_h[i])) {
for (uint32_t i = 0; i < N; i++) {
if (out_h_g2[i] != (inp_h[i] * inp_h[i] * inp_h[i])) {
btestPassed2 = false;
break;
}
@@ -247,55 +239,81 @@ TEST_CASE("Unit_hipGraphAddEventWaitNode_MultGraphOneStrmDependency") {
*/
TEST_CASE("Unit_hipGraphAddEventWaitNode_differentFlags") {
SECTION("flag = hipEventBlockingSync") {
validate_hipGraphAddEventWaitNode_internodedep(0, 1,
hipEventBlockingSync);
validate_hipGraphAddEventWaitNode_internodedep(0, 1, hipEventBlockingSync);
}
SECTION("graph = hipEventDisableTiming") {
validate_hipGraphAddEventWaitNode_internodedep(0, 1,
hipEventDisableTiming);
validate_hipGraphAddEventWaitNode_internodedep(0, 1, hipEventDisableTiming);
}
}
/**
* Scenario 6
* Scenario 6: Positive parameter tests
*/
TEST_CASE("Unit_hipGraphAddEventWaitNode_Negative") {
TEST_CASE("Unit_hipGraphAddEventWaitNode_Positive_Parameters") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipEvent_t event;
HIP_CHECK(hipEventCreate(&event));
hipGraphNode_t eventwait;
SECTION("pGraphNode = nullptr") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(nullptr,
graph, nullptr, 0, event));
hipGraphNode_t dep_node = nullptr;
hipGraphNode_t dep_node2 = nullptr;
HIP_CHECK(hipGraphAddEmptyNode(&dep_node, graph, nullptr, 0));
HIP_CHECK(hipGraphAddEmptyNode(&dep_node2, graph, nullptr, 0));
hipGraphNode_t dep_nodes[] = {dep_node, dep_node2};
size_t numDeps = 0;
SECTION("numDependencies is zero, dependencies is not nullptr") {
HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, dep_nodes, 0, event));
HIP_CHECK(hipGraphNodeGetDependencies(eventwait, nullptr, &numDeps));
REQUIRE(numDeps == 0);
}
SECTION("graph = nullptr") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
nullptr, nullptr, 0, event));
SECTION("numDependencies < dependencies length") {
HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, dep_nodes, 1, event));
HIP_CHECK(hipGraphNodeGetDependencies(eventwait, nullptr, &numDeps));
REQUIRE(numDeps == 1);
}
SECTION("pDependencies = nullptr") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
graph, nullptr, 1, event));
}
SECTION("event = nullptr") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
graph, nullptr, 0, nullptr));
}
SECTION("graph is uninitialized") {
hipGraph_t graph_uninit{};
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
graph_uninit, nullptr, 0, event));
}
SECTION("event is uninitialized") {
hipEvent_t event_uninit{};
REQUIRE(hipErrorInvalidValue == hipGraphAddEventWaitNode(&eventwait,
graph, nullptr, 0, event_uninit));
SECTION("numDependencies == dependencies length") {
HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, dep_nodes, 2, event));
HIP_CHECK(hipGraphNodeGetDependencies(eventwait, nullptr, &numDeps));
REQUIRE(numDeps == 2);
}
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event));
}
/**
* Scenario 7
*/
TEST_CASE("Unit_hipGraphAddEventWaitNode_Negative") {
using namespace std::placeholders;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipEvent_t event;
HIP_CHECK(hipEventCreate(&event));
hipGraphNode_t eventwait;
GraphAddNodeCommonNegativeTests(std::bind(hipGraphAddEventWaitNode, _1, _2, _3, _4, event),
graph);
SECTION("event = nullptr") {
HIP_CHECK_ERROR(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, nullptr),
hipErrorInvalidValue);
}
SECTION("graph is uninitialized") {
hipGraph_t graph_uninit{};
HIP_CHECK_ERROR(hipGraphAddEventWaitNode(&eventwait, graph_uninit, nullptr, 0, event),
hipErrorInvalidValue);
}
SECTION("event is uninitialized") {
hipEvent_t event_uninit{};
HIP_CHECK_ERROR(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event_uninit),
hipErrorInvalidValue);
}
HIP_CHECK(hipGraphDestroy(graph));