AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

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

Sfoglia il codice sorgente 
Change-Id: Id3f569d94d347af2f5e27513fa01c5a1e8e30fd9
This commit is contained in:
Nives Vukovic
2023-11-24 21:42:07 +05:30
committato da Maneesh Gupta
parent 133521b22f
commit 575e4cc93e
8 ha cambiato i file con 472 aggiunte e 441 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc
+101 -72
Vedi File
@@ -33,7 +33,12 @@ Testcase Scenarios :
the graph to create an executable graph. Change the event in the
executable graph to event2. Verify that the event record node still
contains event1.
3) Negative Scenarios
3) Scenario to verify that hipGraphExecEventRecordNodeSetEvent can set event
created on different device. Create an event record node with event1 and add it to graph.
Instantiate the graph to create an executable graph. Call the API to change the event in the
executable graph to event2 which has been created on different device. Verify that graph can be
launched and no error is reported.
4) Negative Scenarios
- Input executable graph is a nullptr.
- Input node is a nullptr.
- Input event to set is a nullptr.
@@ -45,27 +50,26 @@ Testcase Scenarios :
- Input node is a event wait node.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#define GRID_DIM 512
#define BLK_DIM 512
#define LEN (GRID_DIM * BLK_DIM)
/**
* Kernel Functions to copy.
*/
static __global__ void copy_ker_func(int* a, int* b) {
int tx = blockIdx.x*blockDim.x + threadIdx.x;
if (tx < LEN) b[tx] = a[tx];
static __global__ void copy_ker_func(int* a, int* b, size_t N) {
int tx = blockIdx.x * blockDim.x + threadIdx.x;
if (tx < N) b[tx] = a[tx];
}
/**
* Scenario 1: Functional scenario (See description Above)
*/
TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") {
size_t memsize = LEN*sizeof(int);
constexpr size_t gridSize = 512;
constexpr size_t blockSize = 512;
constexpr size_t N = gridSize * blockSize;
size_t memsize = N * sizeof(int);
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
// Create events
@@ -75,10 +79,8 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") {
HIP_CHECK(hipEventCreate(&event2_end));
// Create nodes with event_start and event1_end
hipGraphNode_t event_start_rec, event_end_rec;
HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0,
event_start));
HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0,
event1_end));
HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0, event_start));
HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0, event1_end));
int *inp_h, *inp_d, *out_h, *out_d;
// Allocate host buffers
inp_h = reinterpret_cast<int*>(malloc(memsize));
@@ -89,7 +91,7 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") {
HIP_CHECK(hipMalloc(&inp_d, memsize));
HIP_CHECK(hipMalloc(&out_d, 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[i] = 0;
}
@@ -97,44 +99,39 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") {
// Create memcpy and kernel nodes for graph
hipGraphNode_t memcpyH2D, memcpyD2H, kernelnode;
hipKernelNodeParams kernelNodeParams{};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, inp_d,
inp_h, memsize, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nullptr, 0,
out_h, out_d, memsize, hipMemcpyDeviceToHost));
void* kernelArgs1[] = {&inp_d, &out_d};
kernelNodeParams.func = reinterpret_cast<void *>(copy_ker_func);
kernelNodeParams.gridDim = dim3(GRID_DIM);
kernelNodeParams.blockDim = dim3(BLK_DIM);
size_t NElem{N};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, inp_d, inp_h, memsize,
hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nullptr, 0, out_h, out_d, memsize,
hipMemcpyDeviceToHost));
void* kernelArgs1[] = {&inp_d, &out_d, reinterpret_cast<void*>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void*>(copy_ker_func);
kernelNodeParams.gridDim = dim3(gridSize);
kernelNodeParams.blockDim = dim3(blockSize);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode, graph, nullptr, 0,
&kernelNodeParams));
HIP_CHECK(hipGraphAddKernelNode(&kernelnode, graph, nullptr, 0, &kernelNodeParams));
// Create dependencies for graph
HIP_CHECK(hipGraphAddDependencies(graph, &event_start_rec,
&memcpyH2D, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D,
&kernelnode, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode,
&memcpyD2H, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H,
&event_end_rec, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &event_start_rec, &memcpyH2D, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D, &kernelnode, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode, &memcpyD2H, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H, &event_end_rec, 1));
// Instantiate and launch the graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Change the event at event_end_rec node to event2_end
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec,
event_end_rec, event2_end));
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec, event_end_rec, event2_end));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
// Wait for graph to complete
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Validate output
bool btestPassed = true;
for (uint32_t i = 0; i < LEN; i++) {
for (uint32_t i = 0; i < N; i++) {
if (out_h[i] != inp_h[i]) {
btestPassed = false;
break;
@@ -147,8 +144,7 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") {
REQUIRE(t > 0.0f);
// Since event1_end is never recorded, hipEventElapsedTime
// should return error code.
REQUIRE(hipErrorInvalidResourceHandle ==
hipEventElapsedTime(&t, event_start, event1_end));
HIP_CHECK_ERROR(hipEventElapsedTime(&t, event_start, event1_end), hipErrorInvalidResourceHandle);
// Free resources
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
@@ -173,12 +169,10 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_VerifyEventNotChanged") {
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreate(&event2));
hipGraphNode_t eventrec;
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
event1));
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1));
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec,
eventrec, event2));
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec, event2));
HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out));
// validate set event and get event are same
REQUIRE(event1 == event_out);
@@ -190,7 +184,48 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_VerifyEventNotChanged") {
}
/**
* Scenario 3: Negative Tests
* Scenario 3: This test verifies event in node of the executable graph can be changed to event on
* different device
*/
TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Positive_DifferentDevices") {
const auto device_count = HipTest::getDeviceCount();
if (device_count < 2) {
HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
return;
}
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
hipGraph_t graph;
hipEvent_t event1, event2;
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipEventCreate(&event2));
HIP_CHECK(hipSetDevice(0));
hipGraphNode_t eventrec;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1));
// Verify event on different device can be set in graphExec
// Instantiate and launch the graph
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec, event2));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
// Wait for graph to complete
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Free resources
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event2));
HIP_CHECK(hipEventDestroy(event1))
}
/**
* Scenario 4: Negative Parameter Tests
*/
TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Negative") {
hipGraph_t graph;
@@ -199,11 +234,10 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Negative") {
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreate(&event2));
hipGraphNode_t eventrec;
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
event1));
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1));
// Create memset
constexpr size_t Nbytes = 1024;
char *A_d;
char* A_d;
hipGraphNode_t memset_A;
hipMemsetParams memsetParams{};
HIP_CHECK(hipMalloc(&A_d, Nbytes));
@@ -219,66 +253,61 @@ TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Negative") {
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
SECTION("hGraphExec = nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(nullptr, eventrec, event2));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(nullptr, eventrec, event2),
hipErrorInvalidValue);
}
SECTION("hNode = nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, nullptr, event2));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec, nullptr, event2),
hipErrorInvalidValue);
}
SECTION("event = nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec, nullptr));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec, nullptr),
hipErrorInvalidValue);
}
SECTION("hGraphExec is uninitialized") {
hipGraphExec_t graphExec1{};
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec1, eventrec, event2));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec1, eventrec, event2),
hipErrorInvalidValue);
}
SECTION("hNode is uninitialized") {
hipGraphNode_t dummy{};
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, dummy, event2));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec, dummy, event2),
hipErrorInvalidValue);
}
SECTION("event is uninitialized") {
hipEvent_t event_dummy{};
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec,
event_dummy));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec, event_dummy),
hipErrorInvalidValue);
}
SECTION("event record node does not exist") {
hipGraph_t graph1;
HIP_CHECK(hipGraphCreate(&graph1, 0));
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph1, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph1, nullptr, 0, &memsetParams));
hipGraphExec_t graphExec1;
HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec1, eventrec, event2));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec1, eventrec, event2),
hipErrorInvalidValue);
HIP_CHECK(hipGraphExecDestroy(graphExec1));
HIP_CHECK(hipGraphDestroy(graph1));
}
SECTION("pass memset node as hNode") {
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
&memsetParams));
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, memset_A, event2));
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec, memset_A, event2),
hipErrorInvalidValue);
}
SECTION("pass event wait node as hNode") {
hipGraphNode_t event_wait_node;
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0,
event1));
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, event_wait_node,
event2));
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0, event1));
HIP_CHECK_ERROR(hipGraphExecEventRecordNodeSetEvent(graphExec, event_wait_node, event2),
hipErrorInvalidValue);
}
HIP_CHECK(hipFree(A_d));