diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddEventRecordNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddEventRecordNode.cc index 44b2895cb0..4f5b3b0daa 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddEventRecordNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddEventRecordNode.cc @@ -40,19 +40,26 @@ end. Instantiate and Launch the Graph. Wait for the event to complete. Verify that hipEventElapsedTime() returns error. 6) Validate scenario 2 by running the graph multiple times in a loop (100 times) after instantiation. - 7) Negative Scenarios + 7) Validate that no error is reported when numDeps <= dependencies length + 8) Negative Scenarios - Output node is a nullptr. - Input graph is a nullptr. - Input dependencies is a nullptr. + - Node in dependency is from different graph + - Invalid numNodes + - Duplicate node in dependencies - Input event is a nullptr. - Input graph is uninitialized. - Input event is uninitialized. */ +#include -#include #include +#include #include +#include "graph_tests_common.hh" + /** * Scenario 1: Create s simple graph with just one event record * node and instantiate and launch the graph. @@ -66,8 +73,7 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_Simple") { hipEvent_t event; HIP_CHECK(hipEventCreate(&event)); hipGraphNode_t eventrec; - HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event)); // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); @@ -82,8 +88,8 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_Simple") { /** * Local test function */ -static void validateAddEventRecordNode(bool measureTime, bool withFlags, - int nstep, unsigned flag = 0) { +static void validateAddEventRecordNode(bool measureTime, bool withFlags, int nstep, + unsigned flag = 0) { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); constexpr auto blocksPerCU = 6; // to hide latency @@ -111,8 +117,7 @@ static void validateAddEventRecordNode(bool measureTime, bool withFlags, memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); memset(&memsetParams, 0, sizeof(memsetParams)); memsetParams.dst = reinterpret_cast(B_d); memsetParams.value = 0; @@ -120,38 +125,34 @@ static void validateAddEventRecordNode(bool measureTime, bool withFlags, memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, &memsetParams)); - void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; - kernelNodeParams.func = - reinterpret_cast(HipTest::memsetReverse); + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::memsetReverse); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, &kernelNodeParams)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, - A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, - B_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, - C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); - void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&ker_vecAdd, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&ker_vecAdd, graph, nullptr, 0, &kernelNodeParams)); hipEvent_t eventstart, eventend; if (withFlags) { HIP_CHECK(hipEventCreateWithFlags(&eventstart, flag)); @@ -161,10 +162,8 @@ static void validateAddEventRecordNode(bool measureTime, bool withFlags, HIP_CHECK(hipEventCreate(&eventend)); } hipGraphNode_t event_start, event_final; - HIP_CHECK(hipGraphAddEventRecordNode(&event_start, graph, nullptr, 0, - eventstart)); - HIP_CHECK(hipGraphAddEventRecordNode(&event_final, graph, nullptr, 0, - eventend)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_start, graph, nullptr, 0, eventstart)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_final, graph, nullptr, 0, eventend)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph, &event_start, &memset_A, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &event_start, &memset_B, 1)); @@ -260,7 +259,7 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") { HIP_CHECK(hipEventCreateWithFlags(&event_start, hipEventDisableTiming)); HIP_CHECK(hipEventCreateWithFlags(&event_end, hipEventDisableTiming)); // memset node - char *A_d; + char* A_d; hipGraphNode_t memset_A; hipMemsetParams memsetParams{}; HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -271,14 +270,11 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); hipGraphNode_t event_node_start, event_node_end; - HIP_CHECK(hipGraphAddEventRecordNode(&event_node_start, graph, nullptr, 0, - event_start)); - HIP_CHECK(hipGraphAddEventRecordNode(&event_node_end, graph, nullptr, 0, - event_end)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_node_start, graph, nullptr, 0, event_start)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_node_end, graph, nullptr, 0, event_end)); // Add dependencies between nodes HIP_CHECK(hipGraphAddDependencies(graph, &event_node_start, &memset_A, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &event_node_end, 1)); @@ -290,7 +286,7 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") { // Validate hipEventElapsedTime returns error code because timing is // disabled for start and end event nodes. float t; - REQUIRE(hipSuccess != hipEventElapsedTime(&t, event_start, event_end)); + HIP_CHECK_ERROR(hipEventElapsedTime(&t, event_start, event_end), hipErrorInvalidHandle); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipFree(A_d)); @@ -301,44 +297,73 @@ TEST_CASE("Unit_hipGraphAddEventRecordNode_Functional_TimingDisabled") { } /** - * Scenario 7: All negative tests + * Scenario 7: Positive parameter tests */ -TEST_CASE("Unit_hipGraphAddEventRecordNode_Negative") { +TEST_CASE("Unit_hipGraphAddEventRecordNode_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 == hipGraphAddEventRecordNode(nullptr, - graph, nullptr, 0, event)); + hipGraphNode_t eventrec; + + 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(hipGraphAddEventRecordNode(&eventrec, graph, dep_nodes, 0, event)); + HIP_CHECK(hipGraphNodeGetDependencies(eventrec, nullptr, &numDeps)); + REQUIRE(numDeps == 0); } - SECTION("graph = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait, - nullptr, nullptr, 0, event)); + SECTION("numDependencies < dependencies length") { + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, dep_nodes, 1, event)); + HIP_CHECK(hipGraphNodeGetDependencies(eventrec, nullptr, &numDeps)); + REQUIRE(numDeps == 1); } - SECTION("pDependencies = nullptr and numDependencies != 0") { - REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait, - graph, nullptr, 1, event)); - } - - SECTION("event = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait, - graph, nullptr, 0, nullptr)); - } - - SECTION("graph is uninitialized") { - hipGraph_t graph_uninit{}; - REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait, - graph_uninit, nullptr, 0, nullptr)); - } - - SECTION("event is uninitialized") { - hipEvent_t event_uninit{}; - REQUIRE(hipErrorInvalidValue == hipGraphAddEventRecordNode(&eventwait, - graph, nullptr, 0, event_uninit)); + SECTION("numDependencies == dependencies length") { + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, dep_nodes, 2, event)); + HIP_CHECK(hipGraphNodeGetDependencies(eventrec, nullptr, &numDeps)); + REQUIRE(numDeps == 2); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event)); +} + +/** + * Scenario 8: All negative tests + */ +TEST_CASE("Unit_hipGraphAddEventRecordNode_Negative") { + using namespace std::placeholders; + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + hipGraphNode_t eventrec; + + GraphAddNodeCommonNegativeTests(std::bind(hipGraphAddEventRecordNode, _1, _2, _3, _4, event), + graph); + + SECTION("event = nullptr") { + HIP_CHECK_ERROR(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("graph is uninitialized") { + hipGraph_t graph_uninit{}; + HIP_CHECK_ERROR(hipGraphAddEventRecordNode(&eventrec, graph_uninit, nullptr, 0, event), + hipErrorInvalidValue); + } + + SECTION("event is uninitialized") { + hipEvent_t event_uninit{}; + HIP_CHECK_ERROR(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event_uninit), + hipErrorInvalidValue); } HIP_CHECK(hipGraphDestroy(graph)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddEventWaitNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddEventWaitNode.cc index a196ac7428..89e7e03d7e 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddEventWaitNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddEventWaitNode.cc @@ -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 -#include #include +#include #include -#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(&NElem)}; - kernelNodeParams1.func = - reinterpret_cast(HipTest::vector_square); + void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast(&NElem)}; + kernelNodeParams1.func = reinterpret_cast(HipTest::vector_square); kernelNodeParams1.gridDim = dim3(blocks); kernelNodeParams1.blockDim = dim3(threadsPerBlock); kernelNodeParams1.sharedMemBytes = 0; kernelNodeParams1.kernelParams = reinterpret_cast(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(&NElem)}; - kernelNodeParams2.func = - reinterpret_cast(HipTest::vector_cubic); + void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast(&NElem)}; + kernelNodeParams2.func = reinterpret_cast(HipTest::vector_cubic); kernelNodeParams2.gridDim = dim3(blocks); kernelNodeParams2.blockDim = dim3(threadsPerBlock); kernelNodeParams2.sharedMemBytes = 0; kernelNodeParams2.kernelParams = reinterpret_cast(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)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc index 799841b6cb..49fab00121 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc @@ -26,11 +26,12 @@ with the event set in hipGraphAddEventRecordNode. - Output event is a nullptr. - Input node is an empty node. - Input node is a memset node. + - Input node is event wait node - Input node is an uninitialized node. */ -#include #include +#include #include /** @@ -42,8 +43,7 @@ static void validateEventRecordNodeGetEvent(unsigned flag) { hipEvent_t event, event_out; HIP_CHECK(hipEventCreateWithFlags(&event, flag)); hipGraphNode_t eventrec; - HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event)); HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out)); // validate set event and get event are same REQUIRE(event == event_out); @@ -77,31 +77,32 @@ TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Functional") { TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Negative") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); - hipEvent_t event, event_out; - HIP_CHECK(hipEventCreate(&event)); - hipGraphNode_t eventrec; - HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, - event)); + hipEvent_t event_out; + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventrec, eventwait; + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event2)); + SECTION("node = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeGetEvent(nullptr, - &event_out)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeGetEvent(nullptr, &event_out), hipErrorInvalidValue); } SECTION("event_out = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeGetEvent(eventrec, - nullptr)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeGetEvent(eventrec, nullptr), hipErrorInvalidValue); } SECTION("input node is empty node") { hipGraphNode_t EmptyGraphNode; HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeGetEvent(EmptyGraphNode, &event_out)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeGetEvent(EmptyGraphNode, &event_out), + hipErrorInvalidValue); } SECTION("input node is memset node") { constexpr size_t Nbytes = 1024; - char *A_d; + char* A_d; hipGraphNode_t memset_A; hipMemsetParams memsetParams{}; HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -112,19 +113,21 @@ TEST_CASE("Unit_hipGraphEventRecordNodeGetEvent_Negative") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeGetEvent(memset_A, &event_out)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeGetEvent(memset_A, &event_out), hipErrorInvalidValue); HIP_CHECK(hipFree(A_d)); } + SECTION("input node is event wait node") { + HIP_CHECK_ERROR(hipGraphEventRecordNodeGetEvent(eventwait, &event_out), hipErrorInvalidValue); + } + SECTION("input node is uninitialized node") { hipGraphNode_t node_unit{}; - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeGetEvent(node_unit, &event_out)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeGetEvent(node_unit, &event_out), hipErrorInvalidValue); } HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc index 9c0df129d6..93bb0418ff 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc @@ -30,14 +30,16 @@ Testcase Scenarios : - Input event parameter is nullptr. - Empty node is passed as input node. - Memset node is passed as input node. + - Event wait node is passed as input node. - Input node is an uninitialized node. - Input event is an uninitialized event. */ -#include #include +#include #include + /** * Local Function: Set Get test */ @@ -49,8 +51,7 @@ static void validateEventRecordNodeSetEvent(unsigned flag) { HIP_CHECK(hipEventCreate(&event1)); HIP_CHECK(hipEventCreateWithFlags(&event2, flag)); hipGraphNode_t eventrec; - HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1)); // Set a different event HIP_CHECK(hipGraphEventRecordNodeSetEvent(eventrec, event2)); HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out)); @@ -73,11 +74,9 @@ static void setEventWaitNode() { HIP_CHECK(hipEventCreate(&event1)); HIP_CHECK(hipEventCreate(&event2)); hipGraphNode_t eventwait; - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event1)); // Set a different event eventwait using hipGraphEventRecordNodeSetEvent - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeSetEvent(eventwait, event2)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(eventwait, event2), hipErrorInvalidValue); // Free resources HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipEventDestroy(event1)); @@ -98,13 +97,11 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { HIP_CHECK(hipEventCreateWithFlags(&event2_end, hipEventDisableTiming)); // Create nodes hipGraphNode_t event_start_rec, event_end_rec; - HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0, - event1_start)); - HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0, - event1_end)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0, event1_start)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0, event1_end)); // Create memset node constexpr size_t Nbytes = 1024; - char *A_d; + char* A_d; hipGraphNode_t memset_A; hipMemsetParams memsetParams{}; HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -115,8 +112,7 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); // Create dependencies // event_start_rec --> memset_A --> event_end_rec HIP_CHECK(hipGraphAddDependencies(graph, &event_start_rec, &memset_A, 1)); @@ -132,8 +128,7 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { // Validate by measuring time difference between event_end_rec & // event_start_rec float t = 0.0f; - REQUIRE(hipSuccess == hipEventElapsedTime(&t, event1_start, - event1_end)); + REQUIRE(hipSuccess == hipEventElapsedTime(&t, event1_start, event1_end)); REQUIRE(t > 0.0f); // Change the event property after instantiation HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_start_rec, event2_start)); @@ -145,8 +140,7 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") { // hipErrorInvalidHandle when events are created using // hipEventDisableTiming flag. t = 0.0f; - REQUIRE(hipErrorInvalidHandle == - hipEventElapsedTime(&t, event2_start, event2_end)); + HIP_CHECK_ERROR(hipEventElapsedTime(&t, event2_start, event2_end), hipErrorInvalidHandle); // Free resources HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipStreamDestroy(streamForGraph)); @@ -185,28 +179,24 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_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)); SECTION("node = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeSetEvent(nullptr, - event2)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(nullptr, event2), hipErrorInvalidValue); } SECTION("event_out = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventRecordNodeSetEvent(eventrec, - nullptr)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(eventrec, nullptr), hipErrorInvalidValue); } SECTION("input node is empty node") { hipGraphNode_t EmptyGraphNode; HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeSetEvent(EmptyGraphNode, event2)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(EmptyGraphNode, event2), hipErrorInvalidValue); } SECTION("input node is memset node") { constexpr size_t Nbytes = 1024; - char *A_d; + char* A_d; hipGraphNode_t memset_A; hipMemsetParams memsetParams{}; HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -217,10 +207,8 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_Negative") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeSetEvent(memset_A, event2)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(memset_A, event2), hipErrorInvalidValue); HIP_CHECK(hipFree(A_d)); } @@ -230,14 +218,12 @@ TEST_CASE("Unit_hipGraphEventRecordNodeSetEvent_Negative") { SECTION("input node is uninitialized node") { hipGraphNode_t node_uninit{}; - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeSetEvent(node_uninit, event2)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(node_uninit, event2), hipErrorInvalidValue); } SECTION("input event is uninitialized") { hipEvent_t event_uninit{}; - REQUIRE(hipErrorInvalidValue == - hipGraphEventRecordNodeSetEvent(eventrec, event_uninit)); + HIP_CHECK_ERROR(hipGraphEventRecordNodeSetEvent(eventrec, event_uninit), hipErrorInvalidValue); } HIP_CHECK(hipGraphDestroy(graph)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc index 1dfefdc7cb..74bbf87c87 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc @@ -26,13 +26,15 @@ with the event set in hipGraphAddEventWaitNode. - Output event parameter is passed as nullptr. - Input node parameter is an empty node. - Input node parameter is a memset node. + - Input node parameter is a event record node. - Input node parameter is an uninitialized node. */ -#include #include +#include #include + /** * Local Function */ @@ -42,8 +44,7 @@ static void validateEventWaitNodeGetEvent(unsigned flag) { hipEvent_t event, event_out; HIP_CHECK(hipEventCreateWithFlags(&event, flag)); hipGraphNode_t eventwait; - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event)); HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out)); // validate set event and get event are same REQUIRE(event == event_out); @@ -77,31 +78,32 @@ TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Functional") { TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Negative") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); - hipEvent_t event, event_out; - HIP_CHECK(hipEventCreate(&event)); - hipGraphNode_t eventwait; - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event)); + hipEvent_t event_out; + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventrec, eventwait; + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event2)); + SECTION("node = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeGetEvent(nullptr, - &event_out)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeGetEvent(nullptr, &event_out), hipErrorInvalidValue); } SECTION("event_out = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeGetEvent(eventwait, - nullptr)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeGetEvent(eventwait, nullptr), hipErrorInvalidValue); } SECTION("input node is empty node") { hipGraphNode_t EmptyGraphNode; HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeGetEvent(EmptyGraphNode, &event_out)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeGetEvent(EmptyGraphNode, &event_out), + hipErrorInvalidValue); } SECTION("input node is memset node") { constexpr size_t Nbytes = 1024; - char *A_d; + char* A_d; hipGraphNode_t memset_A; hipMemsetParams memsetParams{}; HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -112,19 +114,21 @@ TEST_CASE("Unit_hipGraphEventWaitNodeGetEvent_Negative") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeGetEvent(memset_A, &event_out)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeGetEvent(memset_A, &event_out), hipErrorInvalidValue); HIP_CHECK(hipFree(A_d)); } + SECTION("input node is event record node") { + HIP_CHECK_ERROR(hipGraphEventWaitNodeGetEvent(eventrec, &event_out), hipErrorInvalidValue); + } + SECTION("input node is uninitialized") { hipGraphNode_t node_uninit{}; - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeGetEvent(node_uninit, &event_out)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeGetEvent(node_uninit, &event_out), hipErrorInvalidValue); } HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc index 8751ffe8ed..6c3132262d 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc @@ -37,11 +37,10 @@ Testcase Scenarios : - Input event is an uninitialized node. */ -#include #include +#include #include -#define LEN 512 /** * Local Function @@ -54,8 +53,7 @@ static void validateEventWaitNodeSetEvent(unsigned flag) { HIP_CHECK(hipEventCreate(&event1)); HIP_CHECK(hipEventCreateWithFlags(&event2, flag)); hipGraphNode_t eventwait; - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event1)); // Set a different event HIP_CHECK(hipGraphEventWaitNodeSetEvent(eventwait, event2)); HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out)); @@ -78,11 +76,9 @@ static void setEventRecordNode() { 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)); // Set a different event eventrec using hipGraphEventWaitNodeSetEvent - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeSetEvent(eventrec, event2)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(eventrec, event2), hipErrorInvalidValue); // Free resources HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipEventDestroy(event1)); @@ -93,11 +89,12 @@ static void setEventRecordNode() { * Scenario 2 */ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") { - size_t memsize = LEN * sizeof(int); + constexpr size_t N = 512; + 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; @@ -123,67 +120,56 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") { 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(&NElem)}; - kernelNodeParams1.func = - reinterpret_cast(HipTest::vector_square); + void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast(&NElem)}; + kernelNodeParams1.func = reinterpret_cast(HipTest::vector_square); kernelNodeParams1.gridDim = dim3(blocks); kernelNodeParams1.blockDim = dim3(threadsPerBlock); kernelNodeParams1.sharedMemBytes = 0; kernelNodeParams1.kernelParams = reinterpret_cast(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(&NElem)}; - kernelNodeParams2.func = - reinterpret_cast(HipTest::vector_cubic); + void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast(&NElem)}; + kernelNodeParams2.func = reinterpret_cast(HipTest::vector_cubic); kernelNodeParams2.gridDim = dim3(blocks); kernelNodeParams2.blockDim = dim3(threadsPerBlock); kernelNodeParams2.sharedMemBytes = 0; kernelNodeParams2.kernelParams = reinterpret_cast(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)); @@ -198,16 +184,16 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") { 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; } @@ -256,28 +242,24 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_Negative") { HIP_CHECK(hipEventCreate(&event1)); HIP_CHECK(hipEventCreate(&event2)); hipGraphNode_t eventwait; - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event1)); SECTION("node = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent( - nullptr, event2)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(nullptr, event2), hipErrorInvalidValue); } SECTION("event = nullptr") { - REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent( - eventwait, nullptr)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(eventwait, nullptr), hipErrorInvalidValue); } SECTION("input node is empty node") { hipGraphNode_t EmptyGraphNode; HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeSetEvent(EmptyGraphNode, event2)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(EmptyGraphNode, event2), hipErrorInvalidValue); } SECTION("input node is memset node") { constexpr size_t Nbytes = 1024; - char *A_d; + char* A_d; hipGraphNode_t memset_A; hipMemsetParams memsetParams{}; HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -288,10 +270,8 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_Negative") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, - &memsetParams)); - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeSetEvent(memset_A, event2)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(memset_A, event2), hipErrorInvalidValue); HIP_CHECK(hipFree(A_d)); } @@ -301,14 +281,12 @@ TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_Negative") { SECTION("input node is uninitialized node") { hipGraphNode_t node_uninit{}; - REQUIRE(hipErrorInvalidValue == - hipGraphEventWaitNodeSetEvent(node_uninit, event2)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(node_uninit, event2), hipErrorInvalidValue); } SECTION("input event is uninitialized") { hipEvent_t event_uninit{}; - REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent( - eventwait, event_uninit)); + HIP_CHECK_ERROR(hipGraphEventWaitNodeSetEvent(eventwait, event_uninit), hipErrorInvalidValue); } HIP_CHECK(hipGraphDestroy(graph)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc index 5a1cbe9997..e820dc1cdd 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc @@ -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 #include +#include #include -#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(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(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(&NElem)}; + kernelNodeParams.func = reinterpret_cast(copy_ker_func); + kernelNodeParams.gridDim = dim3(gridSize); + kernelNodeParams.blockDim = dim3(blockSize); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(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)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc index f8ee63d6f8..529f8df3c1 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc @@ -47,33 +47,30 @@ Testcase Scenarios : - Pass event record node as input node. */ -#include #include +#include #include -#define GRID_DIM 64 -#define BLK_DIM 256 -#define LEN (GRID_DIM * BLK_DIM) -#define DELAY_IN_MS 2000 - /** * Kernel Functions to perform square and introduce delay in device. */ -static __global__ void sqr_ker_func(int* a, int* b, int clockrate) { - int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; - if (tx < LEN) b[tx] = a[tx]*a[tx]; - uint64_t wait_t = DELAY_IN_MS, - start = clock64()/clockrate, cur; - do { cur = clock64()/clockrate - start;}while (cur < wait_t); +static __global__ void sqr_ker_func(int* a, int* b, size_t N, int clockrate, size_t delayMs) { + int tx = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + if (tx < N) b[tx] = a[tx] * a[tx]; + uint64_t wait_t = delayMs, start = clock64() / clockrate, cur; + do { + cur = clock64() / clockrate - start; + } while (cur < wait_t); } -static __global__ void sqr_ker_func_gfx11(int* a, int* b, int clockrate) { +static __global__ void sqr_ker_func_gfx11(int* a, int* b, size_t N, int clockrate, size_t delayMs) { #if HT_AMD - int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; - if (tx < LEN) b[tx] = a[tx]*a[tx]; - uint64_t wait_t = DELAY_IN_MS, - start = wall_clock64()/clockrate, cur; - do { cur = wall_clock64()/clockrate - start;}while (cur < wait_t); + int tx = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + if (tx < N) b[tx] = a[tx] * a[tx]; + uint64_t wait_t = delayMs, start = wall_clock64() / clockrate, cur; + do { + cur = wall_clock64() / clockrate - start; + } while (cur < wait_t); #endif } @@ -81,7 +78,10 @@ static __global__ void sqr_ker_func_gfx11(int* a, int* b, int clockrate) { * Scenario 1: Test to validate setting different events in executable graph. */ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { - size_t memsize = LEN*sizeof(int); + constexpr size_t gridSize = 64; + constexpr size_t blockSize = 256; + constexpr size_t N = gridSize * blockSize; + size_t memsize = N * sizeof(int); hipGraph_t graph1, graph2; HIP_CHECK(hipGraphCreate(&graph1, 0)); HIP_CHECK(hipGraphCreate(&graph2, 0)); @@ -91,8 +91,7 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { HIP_CHECK(hipEventCreate(&event2)); // Create nodes with event_start and event1_end hipGraphNode_t event_rec; - HIP_CHECK(hipGraphAddEventRecordNode(&event_rec, graph1, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventRecordNode(&event_rec, graph1, nullptr, 0, event1)); int *inp_h, *inp_d, *out_h, *out_d; // Allocate host buffers inp_h = reinterpret_cast(malloc(memsize)); @@ -103,7 +102,7 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { 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; } @@ -112,10 +111,12 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { // MemcpyH2D -> kernel1 -> event_rec hipGraphNode_t memcpyH2D, kernelnode1; hipKernelNodeParams kernelNodeParams1{}; - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d, - inp_h, memsize, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d, inp_h, memsize, + hipMemcpyHostToDevice)); // Get device clock rate int clkRate = 0; + size_t NElem{N}; + size_t delayMs{2000}; if (IsGfx11()) { HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeWallClockRate, 0)); } else { @@ -123,29 +124,25 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { } // kernel1 auto sqr_ker_func_used = IsGfx11() ? sqr_ker_func_gfx11 : sqr_ker_func; - void* kernelArgs[] = {&inp_d, &out_d, reinterpret_cast(&clkRate)}; - kernelNodeParams1.func = reinterpret_cast(sqr_ker_func_used); - kernelNodeParams1.gridDim = dim3(GRID_DIM); - kernelNodeParams1.blockDim = dim3(BLK_DIM); + void* kernelArgs[] = {&inp_d, &out_d, reinterpret_cast(&NElem), + reinterpret_cast(&clkRate), reinterpret_cast(&delayMs)}; + kernelNodeParams1.func = reinterpret_cast(sqr_ker_func_used); + kernelNodeParams1.gridDim = dim3(gridSize); + kernelNodeParams1.blockDim = dim3(blockSize); kernelNodeParams1.sharedMemBytes = 0; kernelNodeParams1.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams1.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelnode1, graph1, nullptr, 0, - &kernelNodeParams1)); + HIP_CHECK(hipGraphAddKernelNode(&kernelnode1, graph1, nullptr, 0, &kernelNodeParams1)); // Create dependencies for graph1 - HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D, - &kernelnode1, 1)); - HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode1, - &event_rec, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D, &kernelnode1, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode1, &event_rec, 1)); // graph2 creation ........... // waitnode(event1) -> MemcpyD2H hipGraphNode_t event_wait_node, memcpyD2H; - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph2, nullptr, 0, - out_h, out_d, memsize, hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0, - event1)); - HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node, - &memcpyD2H, 1)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph2, nullptr, 0, out_h, out_d, memsize, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0, event1)); + HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node, &memcpyD2H, 1)); // Instantiate graph1 and graph2 hipStream_t streamForGraph1, streamForGraph2; hipGraphExec_t graphExec1, graphExec2; @@ -160,8 +157,8 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { HIP_CHECK(hipStreamSynchronize(streamForGraph2)); // Validate output bool btestPassed = true; - for (uint32_t i = 0; i < LEN; i++) { - if (out_h[i] != (inp_h[i]*inp_h[i])) { + for (uint32_t i = 0; i < N; i++) { + if (out_h[i] != (inp_h[i] * inp_h[i])) { btestPassed = false; break; } @@ -170,10 +167,8 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { // hipGraphExecEventWaitNodeSetEvent() TEST // Change the event at event_wait_node node to event2 and // the event at event_rec node to event2. - HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec1, - event_rec, event2)); - HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec2, - event_wait_node, event2)); + HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec1, event_rec, event2)); + HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec2, event_wait_node, event2)); // Launch graph1 and graph2 HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1)); HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2)); @@ -181,8 +176,8 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { HIP_CHECK(hipStreamSynchronize(streamForGraph2)); // Validate output btestPassed = true; - for (uint32_t i = 0; i < LEN; i++) { - if (out_h[i] != (inp_h[i]*inp_h[i])) { + for (uint32_t i = 0; i < N; i++) { + if (out_h[i] != (inp_h[i] * inp_h[i])) { btestPassed = false; break; } @@ -214,12 +209,10 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_VerifyEventNotChanged") { HIP_CHECK(hipEventCreate(&event1)); HIP_CHECK(hipEventCreate(&event2)); hipGraphNode_t eventwait; - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, event1)); hipGraphExec_t graphExec; HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec, - eventwait, event2)); + HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, event2)); HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out)); // validate set event and get event are same REQUIRE(event1 == event_out); @@ -240,13 +233,11 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_Negative") { HIP_CHECK(hipEventCreate(&event1)); HIP_CHECK(hipEventCreate(&event2)); hipGraphNode_t eventrec, eventwait; - HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, - event1)); - HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, - event1)); + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, 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)); @@ -262,62 +253,59 @@ TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_Negative") { HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); SECTION("hGraphExec = nullptr") { - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(nullptr, eventwait, event2)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(nullptr, eventwait, event2), + hipErrorInvalidValue); } SECTION("hNode = nullptr") { - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(graphExec, nullptr, event2)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec, nullptr, event2), + hipErrorInvalidValue); } SECTION("event = nullptr") { - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, nullptr)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, nullptr), + hipErrorInvalidValue); } SECTION("hGraphExec is uninitialized") { hipGraphExec_t graphExec1{}; - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2), + hipErrorInvalidValue); } SECTION("hNode is uninitialized") { hipGraphNode_t dummy{}; - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(graphExec, dummy, event2)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec, dummy, event2), + hipErrorInvalidValue); } SECTION("event is uninitialized") { hipEvent_t event_dummy{}; - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, - event_dummy)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, event_dummy), + hipErrorInvalidValue); } SECTION("event wait 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 == - hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, 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 == - hipGraphExecEventWaitNodeSetEvent(graphExec, memset_A, event2)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, &memsetParams)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec, memset_A, event2), + hipErrorInvalidValue); } SECTION("pass event record node as hNode") { - REQUIRE(hipErrorInvalidValue == - hipGraphExecEventWaitNodeSetEvent(graphExec, eventrec, event2)); + HIP_CHECK_ERROR(hipGraphExecEventWaitNodeSetEvent(graphExec, eventrec, event2), + hipErrorInvalidValue); } HIP_CHECK(hipFree(A_d));