From 639ba1e5f39c1f6b2836fd42606c61bf6a80393b Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Thu, 16 Nov 2023 10:53:38 +0000 Subject: [PATCH] EXSWHTEC-172 - Hip graph launching and execution tests #36 Change-Id: I2cfadf6335a582f593191860d5084cb79fa411b3 [ROCm/hip-tests commit: 133521b22f95d577391544a39d5c3f88ff98b5d7] --- .../hip-tests/catch/unit/graph/CMakeLists.txt | 1 + .../catch/unit/graph/hipGraphExecDestroy.cc | 52 +- .../catch/unit/graph/hipGraphExecUpdate.cc | 325 +++++------- .../catch/unit/graph/hipGraphLaunch.cc | 493 ++++-------------- .../catch/unit/graph/hipGraphLaunch_old.cc | 412 +++++++++++++++ .../catch/unit/graph/hipGraphUpload.cc | 22 +- 6 files changed, 719 insertions(+), 586 deletions(-) create mode 100644 projects/hip-tests/catch/unit/graph/hipGraphLaunch_old.cc diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index 148d92bb88..cc6a0d7935 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -103,6 +103,7 @@ set(TEST_SRC hipGraphKernelNodeSetParams.cc hipGraphExecKernelNodeSetParams.cc hipGraphLaunch.cc + hipGraphLaunch_old.cc hipGraphMemcpyNodeSetParams1D.cc hipGraphExecMemcpyNodeSetParamsToSymbol_old.cc hipGraphExecMemcpyNodeSetParamsToSymbol.cc diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc index 7ed2e0d718..debeb32a52 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecDestroy.cc @@ -20,26 +20,51 @@ THE SOFTWARE. #include /** -Negative Testcase Scenarios : -1) Pass hipGraphExecDestroy with nullptr. -2) Pass hipGraphExecDestroy with un-initilze structure. -3) Destroy graph before exec-graph destroyed and verify no crash occurs. -*/ + * @addtogroup hipGraphExecDestroy hipGraphExecDestroy + * @{ + * @ingroup GraphTest + * `hipGraphExecDestroy(hipGraphExec_t graphExec)` - + * Destroys an executable graph + */ + +/** + * Test Description + * ------------------------ + * - Test to verify API behavior with invalid arguments: + * -# GraphExec is nullptr + * -# GraphExec is uninitialized + * Test source + * ------------------------ + * - unit/graph/hipGraphExecDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphExecDestroy_Negative_Parameters") { -TEST_CASE("Unit_hipGraphExecDestroy_Negative") { - hipError_t ret; SECTION("Pass hipGraphExecDestroy with nullptr") { - ret = hipGraphExecDestroy(nullptr); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphExecDestroy(nullptr), hipErrorInvalidValue); } + SECTION("Pass hipGraphExecDestroy with un-initilze structure") { - hipGraphExec_t graphExec{}; - ret = hipGraphExecDestroy(graphExec); - REQUIRE(hipErrorInvalidValue == ret); + hipGraphExec_t graph_exec{}; + HIP_CHECK_ERROR(hipGraphExecDestroy(graph_exec), hipErrorInvalidValue); } } -TEST_CASE("Unit_hipGraphExecDestroy_Sequence") { +/** + * Test Description + * ------------------------ + * - Basic positive test for hipGraphExecDestroy + * - create an executable graph and then destroy it + * Test source + * ------------------------ + * - unit/graph/hipGraphExecDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphExecDestroy_Positive_Basic") { hipGraph_t graph; hipGraphExec_t graphExec; hipStream_t streamForGraph; @@ -70,4 +95,3 @@ TEST_CASE("Unit_hipGraphExecDestroy_Sequence") { HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipStreamDestroy(streamForGraph)); } - diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc index ecd7a19f40..bfc7b2c43e 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecUpdate.cc @@ -27,22 +27,6 @@ THE SOFTWARE. * and perform the update if possible. */ -/** -Testcase Scenarios : -Functional- -1) Make a clone of the created graph and update the executable-graph from a clone or same graph again. -2) Update the executable-graph from a graph and make sure they are taking effect. -Negative- -1) When Pass hGraphExec as nullptr and verify api returns error code. -2) When Pass hGraph as nullptr and verify api returns error code. -3) When Pass hErrorNode_out as nullptr and verify api returns error code. -4) When Pass updateResult_out as nullptr and verify api returns error code. -5) When the a graphExec was updated with with different type of node and verify api returns error code. -6) When a node is deleted in hGraph but not its pair from hGraphExec and verify api returns error code. -7) When a node is deleted in hGraphExec but not its pair from hGraph and verify api returns error code. -8) When grpah dependencies differ but graph have same node and verify api returns error code. -*/ - #include #include #include @@ -65,13 +49,11 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Basic") { hipGraphNode_t hErrorNode_out{}; hipGraphExecUpdateResult updateResult_out{}; SECTION("Pass hGraphExec as nullptr") { - ret = hipGraphExecUpdate(nullptr, graph, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(nullptr, graph, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorInvalidValue == ret); } SECTION("Pass hGraph as nullptr") { - ret = hipGraphExecUpdate(graphExec, nullptr, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec, nullptr, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorInvalidValue == ret); } SECTION("Pass hErrorNode_out as nullptr") { @@ -101,10 +83,9 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(char); constexpr size_t val = 0; - char *devData; + char* devData; int *A_d, *A_h; - HipTest::initArrays(&A_d, nullptr, nullptr, - &A_h, nullptr, nullptr, N, false); + HipTest::initArrays(&A_d, nullptr, nullptr, &A_h, nullptr, nullptr, N, false); HIP_CHECK(hipMalloc(&devData, Nbytes)); hipGraph_t graph, graph2; hipGraphExec_t graphExec; @@ -122,18 +103,16 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams)); std::vector dependencies; dependencies.push_back(memsetNode); HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphCreate(&graph2, 0)); HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); // graphExec was created before memcpyTemp was added to graph. - ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipGraphExecUpdateErrorNodeTypeChanged == updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); HIP_CHECK(hipFree(devData)); @@ -164,7 +143,7 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; size_t NElem{N}; - int *hData = reinterpret_cast(malloc(Nbytes)); + int* hData = reinterpret_cast(malloc(Nbytes)); REQUIRE(hData != nullptr); memset(hData, 0, Nbytes); hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, memcpyTemp; @@ -180,57 +159,52 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph1, 0)); HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); + void* kernelArgs[] = {&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(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph1, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph1, nullptr, 0, &kernelNodeParams)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_A, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_B, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &kernel_vecAdd, &memcpy_C, 1)); // Create a cloned graph and added extra node to it HIP_CHECK(hipGraphClone(&graph2, graph1)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyTemp, graph2, nullptr, 0, - C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyTemp, graph2, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); SECTION("When a node deleted from Graph but not from its pair GraphExec") { - ret = hipGraphExecUpdate(graphExec2, graph1, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec2, graph1, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); } SECTION("When a node deleted from GraphExec but not from its pair Graph") { - ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); } SECTION("When the dependent nodes of a pair differ") { HIP_CHECK(hipGraphCreate(&graph3, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph3, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph3, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph3, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph3, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph3, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph3, nullptr, 0, &kernelNodeParams)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph3, &memcpy_A, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph3, &memcpy_B, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph3, &memcpy_C, &kernel_vecAdd, 1)); - ret = hipGraphExecUpdate(graphExec1, graph3, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec1, graph3, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); HIP_CHECK(hipGraphDestroy(graph3)); } @@ -265,7 +239,7 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; size_t NElem{N}; - int *hData = reinterpret_cast(malloc(Nbytes)); + int* hData = reinterpret_cast(malloc(Nbytes)); REQUIRE(hData != nullptr); memset(hData, 0, Nbytes); hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; @@ -280,22 +254,20 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = - reinterpret_cast(HipTest::vector_square); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vector_square); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSquare, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSquare, graph, nullptr, 0, &kernelNodeParams)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kernel_vecSquare, 1)); HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kernel_vecSquare, 1)); @@ -304,36 +276,32 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); SECTION("Update graphExec with clone graph") { HIP_CHECK(hipGraphClone(&clonedgraph, graph)); - HIP_CHECK(hipGraphExecUpdate(graphExec, clonedgraph, &hErrorNode_out, - &updateResult_out)); + HIP_CHECK(hipGraphExecUpdate(graphExec, clonedgraph, &hErrorNode_out, &updateResult_out)); } // Code for new graph creation with samilar node setup HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphMemcpyNodeSetParams1D(memcpy_C, hData, C_d, Nbytes, - hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphMemcpyNodeSetParams1D(memcpy_C, hData, C_d, Nbytes, hipMemcpyDeviceToHost)); memset(&kernelNodeParams, 0, sizeof(hipKernelNodeParams)); - 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(&kernel_vecAdd, graph2, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph2, nullptr, 0, &kernelNodeParams)); // Create dependencies HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_A, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_B, &kernel_vecAdd, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &kernel_vecAdd, &memcpy_C, 1)); // Update the graphExec graph from graph -> graph2 - HIP_CHECK(hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, - &updateResult_out)); + HIP_CHECK(hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out)); REQUIRE(updateResult_out == hipGraphExecUpdateSuccess); HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); HIP_CHECK(hipStreamSynchronize(streamForGraph)); @@ -380,24 +348,22 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_ParametersChanged") { hipGraphExecUpdateResult updateResult_out; HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); SECTION("Update graphExec with similar graph and verify") { HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out, - &updateResult_out); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipSuccess == ret); HIP_CHECK(hipGraphDestroy(graph2)); } SECTION("Update graphExec with similar graph and verify") { HIP_CHECK(hipGraphCreate(&graph3, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_h, B_d, - Nbytes, hipMemcpyDeviceToHost)); - ret = hipGraphExecUpdate(graphExec1, graph3, &hErrorNode_out, - &updateResult_out); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_h, B_d, Nbytes, + hipMemcpyDeviceToHost)); + ret = hipGraphExecUpdate(graphExec1, graph3, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); REQUIRE(hipGraphExecUpdateErrorParametersChanged == updateResult_out); @@ -437,16 +403,15 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_1") { HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); HIP_CHECK(hipGraphCreate(&graph1, 0)); HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); // When count of nodes directly differ in graphExec1 and graph2 - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); - ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out, - &updateResult_out); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); + ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); #if HT_NVIDIA @@ -495,16 +460,15 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_2") { hipGraphExecUpdateResult updateResult_out; HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); // Delete a node from the graph HIP_CHECK(hipGraphDestroyNode(memcpy_B)); SECTION("When a node deleted from Graph but not from its pair GraphExec") { - ret = hipGraphExecUpdate(graphExec1, graph1, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec1, graph1, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); #if HT_NVIDIA @@ -513,11 +477,10 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_2") { } SECTION("Update the GraphExec with similar graph where a node get deleted") { HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_d, C_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_d, C_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); - ret = hipGraphExecUpdate(graphExec2, graph1, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec2, graph1, &hErrorNode_out, &updateResult_out); #if HT_NVIDIA REQUIRE(hipErrorGraphExecUpdateFailure == ret); REQUIRE(hipGraphExecUpdateErrorNotSupported == updateResult_out); @@ -529,13 +492,12 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_2") { } SECTION("When A node is deleted in GraphExec but not its pair from Graph") { HIP_CHECK(hipGraphCreate(&graph3, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph3, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph3, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphInstantiate(&graphExec3, graph3, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - ret = hipGraphExecUpdate(graphExec3, graph3, &hErrorNode_out, - &updateResult_out); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + ret = hipGraphExecUpdate(graphExec3, graph3, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); #if HT_NVIDIA REQUIRE(hipGraphExecUpdateErrorNotSupported == updateResult_out); @@ -581,27 +543,26 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Dependent_NodesDiffer") { hipGraphExecUpdateResult updateResult_out; HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_d, C_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_d, C_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_A, &memcpy_B, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_B, &memcpy_C, 1)); HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_d, C_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_d, C_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_A, &memcpy_C, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_B, &memcpy_C, 1)); - ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); @@ -642,10 +603,10 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_NodeType_Changed") { HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_A, &memcpy_B, 1)); HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); @@ -658,13 +619,11 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_NodeType_Changed") { memsetParams.elementSize = sizeof(char); memsetParams.width = Nbytes; memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph2, nullptr, 0, - &memsetParams)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph2, nullptr, 0, &memsetParams)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_A, &memsetNode, 1)); - ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); #if HT_NVIDIA REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); @@ -726,22 +685,21 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_MultiDevice_Context_Changed") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&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(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecADD, graph1, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecADD, graph1, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_A, &kernel_vecADD, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_B, &kernel_vecADD, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &kernel_vecADD, &memcpy_C, 1)); @@ -750,27 +708,25 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_MultiDevice_Context_Changed") { HIP_CHECK(hipSetDevice(1)); HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); memset(&kernelNodeParams, 0x00, sizeof(hipKernelNodeParams)); - void* kernelArgs1[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorSUB); + void* kernelArgs1[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorSUB); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSUB, graph2, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSUB, graph2, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_A, &kernel_vecSUB, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_B, &kernel_vecSUB, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &kernel_vecSUB, &memcpy_C, 1)); - ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); REQUIRE(hipGraphExecUpdateErrorUnsupportedFunctionChange == updateResult_out); @@ -819,49 +775,46 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional_KernelFunction_Changed") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipGraphCreate(&graph1, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); hipKernelNodeParams kernelNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&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(kernelArgs); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecADD, graph1, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecADD, graph1, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_A, &kernel_vecADD, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_B, &kernel_vecADD, 1)); HIP_CHECK(hipGraphAddDependencies(graph1, &kernel_vecADD, &memcpy_C, 1)); HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); HIP_CHECK(hipGraphCreate(&graph2, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); memset(&kernelNodeParams, 0x00, sizeof(hipKernelNodeParams)); - void* kernelArgs1[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorSUB); + void* kernelArgs1[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorSUB); kernelNodeParams.gridDim = dim3(blocks); kernelNodeParams.blockDim = dim3(threadsPerBlock); kernelNodeParams.sharedMemBytes = 0; kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSUB, graph2, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSUB, graph2, nullptr, 0, &kernelNodeParams)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_A, &kernel_vecSUB, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_B, &kernel_vecSUB, 1)); HIP_CHECK(hipGraphAddDependencies(graph2, &kernel_vecSUB, &memcpy_C, 1)); - ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, - &updateResult_out); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out); REQUIRE(hipSuccess == ret); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); diff --git a/projects/hip-tests/catch/unit/graph/hipGraphLaunch.cc b/projects/hip-tests/catch/unit/graph/hipGraphLaunch.cc index 21afeffce8..e0513a28e0 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphLaunch.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphLaunch.cc @@ -19,394 +19,127 @@ THE SOFTWARE. #include #include -#include -/* Test verifies hipGraphLaunch API -Negative scenarios - -1) Pass graphExec as nullptr and verify api returns error code. -2) Pass pGraphExec as nullptr and stream as hipStreamPerThread and verify api returns error code. -3) Pass pGraphExec as empty object and verify api returns error code. -4) Destroy executable graph and try to launch it. Make sure api should not crash and it should returns error code. -5) Destroy stream and try to launch respective executable graph. Make sure api should not crash and it should returns error code. -6) Destroy actual graph created and try to launch respective executable graph. - Check api should execute properly without crash or error code. -Functional Scenario - -1) Check basic functionality with stream as hipStreamPerThread -2) Test hipGraphLaunch call on multiple devices. -3) Create a graph with multiple nodes. Create an executable graph. - Launch the executable graph 3 times in stream simultaneously. - Wait for stream. Validate the output. No issues should be observed -4) Create a graph with multiple nodes. Create an executable graph. - Verify if an executable graph be launched on null stream. -*/ -#define SIZE 1024 -#define TEST_LOOP_SIZE 3 +/** + * @addtogroup hipGraphLaunch hipGraphLaunch + * @{ + * @ingroup GraphTest + * `hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream)` - + * Launches an executable graph in a stream + */ -TEST_CASE("Unit_hipGraphLaunch_Negative") { - hipError_t ret; - SECTION("Pass pGraphExec as nullptr") { - hipStream_t stream{}; +static void HostFunctionSetToZero(void* arg) { + int* test_number = (int*)arg; + (*test_number) = 0; +} + +static void HostFunctionAddOne(void* arg) { + int* test_number = (int*)arg; + (*test_number) += 1; +} + +/* create an executable graph that will set an integer pointed to by 'number' to one*/ +static void CreateTestExecutableGraph(hipGraphExec_t* graph_exec, int* number) { + hipGraph_t graph; + hipGraphNode_t node_error; + + hipGraphNode_t node_set_zero; + hipHostNodeParams params_set_to_zero = {HostFunctionSetToZero, number}; + + hipGraphNode_t node_add_one; + hipHostNodeParams params_set_add_one = {HostFunctionAddOne, number}; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HIP_CHECK(hipGraphAddHostNode(&node_set_zero, graph, nullptr, 0, ¶ms_set_to_zero)); + HIP_CHECK(hipGraphAddHostNode(&node_add_one, graph, &node_set_zero, 1, ¶ms_set_add_one)); + + HIP_CHECK(hipGraphInstantiate(graph_exec, graph, &node_error, nullptr, 0)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +static void HipGraphLaunch_Positive_Simple(hipStream_t stream) { + int number = 5; + + hipGraphExec_t graph_exec; + CreateTestExecutableGraph(&graph_exec, &number); + + HIP_CHECK(hipGraphLaunch(graph_exec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(number == 1); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); +} + + +/** + * Test Description + * ------------------------ + * - Basic positive test for hipGraphLaunch + * -# stream as a created stream + * -# with stream as hipStreamPerThread + * Test source + * ------------------------ + * - unit/graph/hipGraphLaunch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphLaunch_Positive") { + SECTION("stream as a created stream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HipGraphLaunch_Positive_Simple(stream); + HIP_CHECK(hipStreamDestroy(stream)); + } + + SECTION("with stream as hipStreamPerThread") { + HipGraphLaunch_Positive_Simple(hipStreamPerThread); + } +} + +/** + * Test Description + * ------------------------ + * - Negative parameter test for hipGraphLaunch + * -# graphExec is nullptr and stream is a created stream + * -# graphExec is nullptr and stream is hipStreamPerThread + * -# graphExec is an empty object + * -# graphExec is destroyed before calling hipGraphLaunch + * Test source + * ------------------------ + * - unit/graph/hipGraphLaunch.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphLaunch_Negative_Parameters") { + SECTION("graphExec is nullptr and stream is a created stream") { + hipStream_t stream; + hipError_t ret; + HIP_CHECK(hipStreamCreate(&stream)); ret = hipGraphLaunch(nullptr, stream); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pGraphExec as nullptr and stream as hipStreamPerThread") { - ret = hipGraphLaunch(nullptr, hipStreamPerThread); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pGraphExec as empty object") { - hipGraphExec_t graphExec{}; - hipStream_t stream{}; - ret = hipGraphLaunch(graphExec, stream); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Destroy executable graph and try to launch it") { - constexpr size_t Nbytes = 1024; - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t stream; - hipGraphNode_t memsetNode; - - char *devData; - HIP_CHECK(hipMalloc(&devData, Nbytes)); - - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&stream)); - - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = 0; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - - HIP_CHECK(hipGraphExecDestroy(graphExec)); - // Launch again after destroy graph exec object. - ret = hipGraphLaunch(graphExec, stream); - REQUIRE(hipErrorInvalidValue == ret); - - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphDestroy(graph)); HIP_CHECK(hipStreamDestroy(stream)); + REQUIRE(ret == hipErrorInvalidValue); } -/* In this case in CUDA setup this api call is giving - unknown error (999) - So enabling this test for both AMD and CUDA by checking with hipSuccess */ - SECTION("Destroy stream and try to launch respective executable graph") { - constexpr size_t Nbytes = 1024; - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t stream; - hipGraphNode_t memsetNode; - char *devData; - HIP_CHECK(hipMalloc(&devData, Nbytes)); - - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&stream)); - - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = 0; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - - HIP_CHECK(hipStreamDestroy(stream)); - // Launch again after destroy stream - ret = hipGraphLaunch(graphExec, stream); - REQUIRE(hipSuccess != ret); - - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipGraphDestroy(graph)); + SECTION("graphExec is nullptr and stream is hipStreamPerThread") { + HIP_CHECK_ERROR(hipGraphLaunch(nullptr, hipStreamPerThread), hipErrorInvalidValue); } - SECTION("Destroy graph and try to launch respective executable graph") { - constexpr size_t Nbytes = 1024; - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t stream; - hipGraphNode_t memsetNode; - char *devData; - HIP_CHECK(hipMalloc(&devData, Nbytes)); + SECTION("graphExec is an empty object") { + hipGraphExec_t graph_exec{}; + HIP_CHECK_ERROR(hipGraphLaunch(graph_exec, hipStreamPerThread), hipErrorInvalidValue); + } - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&stream)); - - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(devData); - memsetParams.value = 0; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - - HIP_CHECK(hipGraphDestroy(graph)); - // Launch again after destroy graph - ret = hipGraphLaunch(graphExec, stream); - REQUIRE(hipSuccess == ret); - - HIP_CHECK(hipFree(devData)); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipStreamDestroy(stream)); + SECTION("graphExec is destroyed") { + int number = 5; + hipGraphExec_t graph_exec; + CreateTestExecutableGraph(&graph_exec, &number); + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + REQUIRE(number == 1); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK_ERROR(hipGraphLaunch(graph_exec, hipStreamPerThread), hipErrorInvalidValue); } } - -TEST_CASE("Unit_hipGraphLaunch_Functional_hipStreamPerThread") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(char); - constexpr size_t val = 0; - constexpr size_t updateVal = 2; - char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - char *A_h{nullptr}, *B_h{nullptr}; - - HipTest::initArrays(&A_d, &B_d, &C_d, - &A_h, &B_h, nullptr, N, false); - - hipGraph_t graph; - hipGraphExec_t graphExec; - hipGraphNode_t memsetNode; - - HIP_CHECK(hipGraphCreate(&graph, 0)); - - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(C_d); - memsetParams.value = val; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - - std::vector dependencies; - dependencies.push_back(memsetNode); - - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(A_d); - memsetParams.value = updateVal; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), - dependencies.size(), &memsetParams)); - HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); - dependencies.push_back(memsetNode); - - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, hipStreamPerThread)); - HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - - HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); - - // Validating the result - for (size_t i = 0; i < Nbytes; i++) { - if (A_h[i] != updateVal) { - WARN("Validation failed at- " << i << " A_h[i] " << A_h[i]); - REQUIRE(false); - } - } - - HipTest::freeArrays(A_d, B_d, C_d, - A_h, B_h, nullptr, false); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipGraphDestroy(graph)); -} - -static void hipGraphLaunch_test() { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(char); - constexpr size_t val = 0; - constexpr size_t updateVal = 1; - char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - char *A_h{nullptr}, *B_h{nullptr}; - - HipTest::initArrays(&A_d, &B_d, &C_d, - &A_h, &B_h, nullptr, N, false); - - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t streamForGraph; - hipGraphNode_t memsetNode; - - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); - - hipMemsetParams memsetParams{}; - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(C_d); - memsetParams.value = val; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, - &memsetParams)); - - std::vector dependencies; - dependencies.push_back(memsetNode); - - memset(&memsetParams, 0, sizeof(memsetParams)); - memsetParams.dst = reinterpret_cast(A_d); - memsetParams.value = updateVal; - memsetParams.pitch = 0; - memsetParams.elementSize = sizeof(char); - memsetParams.width = Nbytes; - memsetParams.height = 1; - HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), - dependencies.size(), &memsetParams)); - HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); - dependencies.push_back(memsetNode); - - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); - - HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); - - // Validating the result - for (size_t i = 0; i < Nbytes; i++) { - if (A_h[i] != updateVal) { - WARN("Validation failed at- " << i << " A_h[i] " << A_h[i]); - REQUIRE(false); - } - } - - HipTest::freeArrays(A_d, B_d, C_d, - A_h, B_h, nullptr, false); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); -} - -TEST_CASE("Unit_hipGraphLaunch_Functional_multidevice_test") { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - - if (numDevices > 0) { - for (int i = 0; i < numDevices; i++) { - HIP_CHECK(hipSetDevice(i)); - hipGraphLaunch_test(); - } - } else { - SUCCEED("Skipped the testcase as there is no device to test."); - } -} - -// Function to fill input data -static void fillRandInpData(int *A1_h, int *A2_h, size_t N) { - unsigned int seed = time(nullptr); - for (size_t i = 0; i < N; i++) { - A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF); - A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF); - } -} -// Function to validate result -static void validateOutData(int *A1_h, int *A2_h, size_t N) { - for (size_t i = 0; i < N; i++) { - int result = (A1_h[i]*A1_h[i]); - REQUIRE(result == A2_h[i]); - } -} -/* - * 1.Create a graph with multiple nodes. Create an executable graph. - * Launch the executable graph 3 times in stream simultaneously. - * Wait for stream. Validate the output. No issues should be observed - * 2.Create a graph with multiple nodes. Create an executable graph. - * Verify if an executable graph be launched on null stream. -*/ -TEST_CASE("Unit_hipGraphLaunch_Functional_MultipleLaunch") { - size_t memSize = SIZE; - constexpr auto blocksPerCU = 6; // to hide latency - constexpr auto threadsPerBlock = 256; - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, - threadsPerBlock, SIZE); - hipGraph_t graph; - std::vector nodeDependencies; - - HIP_CHECK(hipGraphCreate(&graph, 0)); - int *A_h{nullptr}, *A_d{nullptr}, *C_d{nullptr}, *C_h{nullptr}; - - HipTest::initArrays(&A_d, &C_d, nullptr, - &A_h, &C_h, nullptr, SIZE, false); - - hipGraphNode_t memcpyH2D, memcpyD2H, kernelNode; - - // Create memcpy H2D nodes - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, - 0, A_d, A_h, (sizeof(int)*SIZE), hipMemcpyHostToDevice)); - nodeDependencies.push_back(memcpyH2D); - // Creating kernel node - hipKernelNodeParams kerNodeParams; - void* kernelArgs[] = {reinterpret_cast(&A_d), - reinterpret_cast(&C_d), - reinterpret_cast(&memSize)}; - kerNodeParams.func = reinterpret_cast(HipTest::vector_square); - kerNodeParams.gridDim = dim3(blocks); - kerNodeParams.blockDim = dim3(threadsPerBlock); - kerNodeParams.sharedMemBytes = 0; - kerNodeParams.kernelParams = reinterpret_cast(kernelArgs); - kerNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(), - nodeDependencies.size(), &kerNodeParams)); - nodeDependencies.clear(); - nodeDependencies.push_back(kernelNode); - - // Create memcpy D2H nodes - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nodeDependencies.data(), - nodeDependencies.size(), C_h, C_d, (sizeof(int)*SIZE), - hipMemcpyDeviceToHost)); - nodeDependencies.clear(); - - // Create executable graph - hipStream_t streamForGraph; - hipGraphExec_t graphExec{nullptr}; - HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, - nullptr, 0)); - // Execute graph - SECTION("Multiple Graph Launch") { - for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { - fillRandInpData(A_h, C_h, SIZE); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); - validateOutData(A_h, C_h, SIZE); - } - } - SECTION("Graph launch on Null stream") { - for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { - fillRandInpData(A_h, C_h, SIZE); - HIP_CHECK(hipGraphLaunch(graphExec, 0)); - HIP_CHECK(hipStreamSynchronize(0)); - validateOutData(A_h, C_h, SIZE); - } - } - - HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); - - // Free - HipTest::freeArrays(A_d, C_d, nullptr, A_h, C_h, nullptr, false); -} diff --git a/projects/hip-tests/catch/unit/graph/hipGraphLaunch_old.cc b/projects/hip-tests/catch/unit/graph/hipGraphLaunch_old.cc new file mode 100644 index 0000000000..21afeffce8 --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipGraphLaunch_old.cc @@ -0,0 +1,412 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +/* Test verifies hipGraphLaunch API +Negative scenarios - +1) Pass graphExec as nullptr and verify api returns error code. +2) Pass pGraphExec as nullptr and stream as hipStreamPerThread and verify api returns error code. +3) Pass pGraphExec as empty object and verify api returns error code. +4) Destroy executable graph and try to launch it. Make sure api should not crash and it should returns error code. +5) Destroy stream and try to launch respective executable graph. Make sure api should not crash and it should returns error code. +6) Destroy actual graph created and try to launch respective executable graph. + Check api should execute properly without crash or error code. +Functional Scenario - +1) Check basic functionality with stream as hipStreamPerThread +2) Test hipGraphLaunch call on multiple devices. +3) Create a graph with multiple nodes. Create an executable graph. + Launch the executable graph 3 times in stream simultaneously. + Wait for stream. Validate the output. No issues should be observed +4) Create a graph with multiple nodes. Create an executable graph. + Verify if an executable graph be launched on null stream. +*/ + +#define SIZE 1024 +#define TEST_LOOP_SIZE 3 + +TEST_CASE("Unit_hipGraphLaunch_Negative") { + hipError_t ret; + SECTION("Pass pGraphExec as nullptr") { + hipStream_t stream{}; + ret = hipGraphLaunch(nullptr, stream); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pGraphExec as nullptr and stream as hipStreamPerThread") { + ret = hipGraphLaunch(nullptr, hipStreamPerThread); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pGraphExec as empty object") { + hipGraphExec_t graphExec{}; + hipStream_t stream{}; + ret = hipGraphLaunch(graphExec, stream); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Destroy executable graph and try to launch it") { + constexpr size_t Nbytes = 1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t memsetNode; + + char *devData; + HIP_CHECK(hipMalloc(&devData, Nbytes)); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + // Launch again after destroy graph exec object. + ret = hipGraphLaunch(graphExec, stream); + REQUIRE(hipErrorInvalidValue == ret); + + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + } +/* In this case in CUDA setup this api call is giving - unknown error (999) + So enabling this test for both AMD and CUDA by checking with hipSuccess */ + SECTION("Destroy stream and try to launch respective executable graph") { + constexpr size_t Nbytes = 1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t memsetNode; + + char *devData; + HIP_CHECK(hipMalloc(&devData, Nbytes)); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipStreamDestroy(stream)); + // Launch again after destroy stream + ret = hipGraphLaunch(graphExec, stream); + REQUIRE(hipSuccess != ret); + + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + } + SECTION("Destroy graph and try to launch respective executable graph") { + constexpr size_t Nbytes = 1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t memsetNode; + + char *devData; + HIP_CHECK(hipMalloc(&devData, Nbytes)); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphDestroy(graph)); + // Launch again after destroy graph + ret = hipGraphLaunch(graphExec, stream); + REQUIRE(hipSuccess == ret); + + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +TEST_CASE("Unit_hipGraphLaunch_Functional_hipStreamPerThread") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + constexpr size_t updateVal = 2; + char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + char *A_h{nullptr}, *B_h{nullptr}; + + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, nullptr, N, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(C_d); + memsetParams.value = val; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = updateVal; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), + dependencies.size(), &memsetParams)); + HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + + // Validating the result + for (size_t i = 0; i < Nbytes; i++) { + if (A_h[i] != updateVal) { + WARN("Validation failed at- " << i << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, C_d, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +static void hipGraphLaunch_test() { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + constexpr size_t updateVal = 1; + char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + char *A_h{nullptr}, *B_h{nullptr}; + + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, nullptr, N, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(C_d); + memsetParams.value = val; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = updateVal; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(), + dependencies.size(), &memsetParams)); + HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams)); + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + + // Validating the result + for (size_t i = 0; i < Nbytes; i++) { + if (A_h[i] != updateVal) { + WARN("Validation failed at- " << i << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, C_d, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +TEST_CASE("Unit_hipGraphLaunch_Functional_multidevice_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices > 0) { + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipSetDevice(i)); + hipGraphLaunch_test(); + } + } else { + SUCCEED("Skipped the testcase as there is no device to test."); + } +} + +// Function to fill input data +static void fillRandInpData(int *A1_h, int *A2_h, size_t N) { + unsigned int seed = time(nullptr); + for (size_t i = 0; i < N; i++) { + A1_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + A2_h[i] = (HipTest::RAND_R(&seed) & 0xFF); + } +} +// Function to validate result +static void validateOutData(int *A1_h, int *A2_h, size_t N) { + for (size_t i = 0; i < N; i++) { + int result = (A1_h[i]*A1_h[i]); + REQUIRE(result == A2_h[i]); + } +} +/* + * 1.Create a graph with multiple nodes. Create an executable graph. + * Launch the executable graph 3 times in stream simultaneously. + * Wait for stream. Validate the output. No issues should be observed + * 2.Create a graph with multiple nodes. Create an executable graph. + * Verify if an executable graph be launched on null stream. +*/ +TEST_CASE("Unit_hipGraphLaunch_Functional_MultipleLaunch") { + size_t memSize = SIZE; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, SIZE); + hipGraph_t graph; + std::vector nodeDependencies; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + int *A_h{nullptr}, *A_d{nullptr}, *C_d{nullptr}, *C_h{nullptr}; + + HipTest::initArrays(&A_d, &C_d, nullptr, + &A_h, &C_h, nullptr, SIZE, false); + + hipGraphNode_t memcpyH2D, memcpyD2H, kernelNode; + + // Create memcpy H2D nodes + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, + 0, A_d, A_h, (sizeof(int)*SIZE), hipMemcpyHostToDevice)); + nodeDependencies.push_back(memcpyH2D); + // Creating kernel node + hipKernelNodeParams kerNodeParams; + void* kernelArgs[] = {reinterpret_cast(&A_d), + reinterpret_cast(&C_d), + reinterpret_cast(&memSize)}; + kerNodeParams.func = reinterpret_cast(HipTest::vector_square); + kerNodeParams.gridDim = dim3(blocks); + kerNodeParams.blockDim = dim3(threadsPerBlock); + kerNodeParams.sharedMemBytes = 0; + kerNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kerNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &kerNodeParams)); + nodeDependencies.clear(); + nodeDependencies.push_back(kernelNode); + + // Create memcpy D2H nodes + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nodeDependencies.data(), + nodeDependencies.size(), C_h, C_d, (sizeof(int)*SIZE), + hipMemcpyDeviceToHost)); + nodeDependencies.clear(); + + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec{nullptr}; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, + nullptr, 0)); + // Execute graph + SECTION("Multiple Graph Launch") { + for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { + fillRandInpData(A_h, C_h, SIZE); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + validateOutData(A_h, C_h, SIZE); + } + } + SECTION("Graph launch on Null stream") { + for (int iter = 0; iter < TEST_LOOP_SIZE; iter++) { + fillRandInpData(A_h, C_h, SIZE); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + validateOutData(A_h, C_h, SIZE); + } + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + + // Free + HipTest::freeArrays(A_d, C_d, nullptr, A_h, C_h, nullptr, false); +} diff --git a/projects/hip-tests/catch/unit/graph/hipGraphUpload.cc b/projects/hip-tests/catch/unit/graph/hipGraphUpload.cc index 53b22e7e6f..f4db37c05f 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphUpload.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphUpload.cc @@ -261,9 +261,10 @@ TEST_CASE("Unit_hipGraphUpload_Functional_With_Priority_Stream") { 1) Pass graphExec node as nullptr. 2) Pass graphExec node as uninitialize object 3) Pass stream as uninitialize object +4) Graphexec is destroyed before upload */ -TEST_CASE("Unit_hipGraphUpload_Negative_Argument_Check") { +TEST_CASE("Unit_hipGraphUpload_Negative_Parameters") { hipGraphExec_t graphExec{}; hipError_t ret; @@ -271,21 +272,30 @@ TEST_CASE("Unit_hipGraphUpload_Negative_Argument_Check") { HIP_CHECK(hipStreamCreate(&stream)); SECTION("Pass graphExec node as nullptr") { - ret = hipGraphUpload(nullptr, stream); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphUpload(nullptr, stream), hipErrorInvalidValue); } SECTION("Pass graphExec node as uninitialize object") { - ret = hipGraphUpload(graphExec, stream); - REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK_ERROR(hipGraphUpload(graphExec, stream), hipErrorInvalidValue); } SECTION("Pass stream as uninitialize object") { hipStream_t stream1{}; hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); ret = hipGraphUpload(graphExec, stream1); REQUIRE(hipSuccess == ret); } + SECTION("graphExec is destroyed"){ + hipGraphExec_t graph_exec; + hipGraph_t graph; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphUpload(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK_ERROR(hipGraphUpload(graph_exec, hipStreamPerThread), hipErrorInvalidValue); + } HIP_CHECK(hipStreamDestroy(stream)); }