diff --git a/catch/unit/graph/hipGraphExecUpdate.cc b/catch/unit/graph/hipGraphExecUpdate.cc index 56c0051802..ecd7a19f40 100644 --- a/catch/unit/graph/hipGraphExecUpdate.cc +++ b/catch/unit/graph/hipGraphExecUpdate.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022-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 @@ -17,6 +17,16 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +/** + * @addtogroup hipGraphExecUpdate hipGraphExecUpdate + * @{ + * @ingroup GraphTest + * `hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, + * hipGraphExecUpdateResultInfo* resultInfo)` - + * Check whether an executable graph can be updated with a graph + * and perform the update if possible. + */ + /** Testcase Scenarios : Functional- @@ -37,7 +47,16 @@ Negative- #include #include -/* Test verifies hipGraphExecUpdate API Negative nullptr check scenarios. +/** + * Test Description + * ------------------------ + * - Test verifies hipGraphExecUpdate API Negative nullptr check scenarios. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 */ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Basic") { hipError_t ret; @@ -45,7 +64,6 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Basic") { hipGraphExec_t graphExec{}; hipGraphNode_t hErrorNode_out{}; hipGraphExecUpdateResult updateResult_out{}; - SECTION("Pass hGraphExec as nullptr") { ret = hipGraphExecUpdate(nullptr, graph, &hErrorNode_out, &updateResult_out); @@ -66,31 +84,36 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_Basic") { } } -/* Test verifies hipGraphExecUpdate API Negative scenarios. - When the a graphExec was updated with with different type of node +/** + * Test Description + * ------------------------ + * - Test verifies hipGraphExecUpdate API Negative scenarios. + * When the a graphExec was updated with with different type of node + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 */ + 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; int *A_d, *A_h; - HipTest::initArrays(&A_d, nullptr, nullptr, &A_h, nullptr, nullptr, N, false); - HIP_CHECK(hipMalloc(&devData, Nbytes)); - hipGraph_t graph, graph2; hipGraphExec_t graphExec; hipStream_t streamForGraph; hipGraphNode_t memsetNode, memcpy_A, hErrorNode_out; hipError_t ret; hipGraphExecUpdateResult updateResult_out; - HIP_CHECK(hipGraphCreate(&graph, 0)); HIP_CHECK(hipStreamCreate(&streamForGraph)); - hipMemsetParams memsetParams{}; memset(&memsetParams, 0, sizeof(memsetParams)); memsetParams.dst = reinterpret_cast(devData); @@ -101,25 +124,18 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") { memsetParams.height = 1; 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)); - // graphExec was created before memcpyTemp was added to graph. ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, &updateResult_out); - REQUIRE(hipGraphExecUpdateErrorNodeTypeChanged == updateResult_out); REQUIRE(hipErrorGraphExecUpdateFailure == ret); - - HIP_CHECK(hipFree(devData)); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); @@ -127,9 +143,19 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") { HIP_CHECK(hipStreamDestroy(streamForGraph)); } -/* Test verifies hipGraphExecUpdate API Negative scenarios. - When the count of nodes differ in hGraphExec and hGraph +/** + * Test Description + * ------------------------ + * - Test verifies hipGraphExecUpdate API Negative scenarios. + * When the count of nodes differ in hGraphExec and hGraph + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 */ + TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { constexpr size_t N = 1024; constexpr size_t Nbytes = N * sizeof(int); @@ -138,11 +164,9 @@ 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)); REQUIRE(hData != nullptr); memset(hData, 0, Nbytes); - hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, memcpyTemp; hipGraphNode_t kernel_vecAdd; hipKernelNodeParams kernelNodeParams{}; @@ -152,10 +176,8 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { hipStream_t streamForGraph; hipGraphNode_t hErrorNode_out; hipGraphExecUpdateResult updateResult_out; - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); 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, @@ -164,7 +186,6 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { 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); @@ -174,20 +195,16 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { kernelNodeParams.extra = nullptr; 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(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); @@ -206,20 +223,17 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { 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); REQUIRE(hipErrorGraphExecUpdateFailure == ret); HIP_CHECK(hipGraphDestroy(graph3)); } - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIP_CHECK(hipGraphExecDestroy(graphExec1)); HIP_CHECK(hipGraphExecDestroy(graphExec2)); @@ -229,10 +243,19 @@ TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") { free(hData); } -/* Functional Scenario - -1) Make a clone of the created graph and update the executable-graph from a clone graph. -2) Update the executable-graph from a graph and make sure they are taking effect. -*/ +/** + * Test Description + * ------------------------ + * - Functional Scenario - + 1) Make a clone of the created graph and update the executable-graph from a clone graph. + 2) Update the executable-graph from a graph and make sure they are taking effect. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { constexpr size_t N = 1024; @@ -242,11 +265,9 @@ 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)); REQUIRE(hData != nullptr); memset(hData, 0, Nbytes); - hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; hipGraphNode_t kernel_vecAdd, kernel_vecSquare; hipKernelNodeParams kernelNodeParams{}; @@ -255,10 +276,8 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { hipStream_t streamForGraph; hipGraphNode_t hErrorNode_out; hipGraphExecUpdateResult updateResult_out; - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); 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, @@ -267,7 +286,6 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { 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); @@ -278,21 +296,17 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { kernelNodeParams.extra = nullptr; 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)); HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecSquare, &memcpy_C, 1)); - // Instantiate and launch the graph 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)); } - // 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, @@ -303,7 +317,6 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { 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); @@ -314,23 +327,18 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { kernelNodeParams.extra = nullptr; 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)); REQUIRE(updateResult_out == hipGraphExecUpdateSuccess); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); HIP_CHECK(hipStreamSynchronize(streamForGraph)); - // Verify graph execution result HipTest::checkVectorADD(A_h, B_h, hData, N); - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipStreamDestroy(streamForGraph)); @@ -339,3 +347,530 @@ TEST_CASE("Unit_hipGraphExecUpdate_Functional") { HIP_CHECK(hipGraphDestroy(clonedgraph)); free(hData); } + +/** + * Test Description + * ------------------------ + * - Functional Basic Check Scenario - 1 + Create a graph1 with memcpy1D node with direction as hipMemcpyHostToDevice + Create a graph2 with memcpy1D node with direction as hipMemcpyHostToDevice + Update graphExec1 with graph2 and verify. It should not return error. + - Negative Scenario - 2 + Create a graph1 with memcpy1D node with direction as hipMemcpyHostToDevice + Instantiate graph1 in graphExec1 + Create a graph2 with memcpy1D node with direction as hipMemcpyDeviceToHost + Update graphExec1 with graph2 and verify. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_ParametersChanged") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphNode_t memcpy_A, memcpy_B; + hipError_t ret; + hipGraph_t graph1, graph2, graph3; + hipGraphExec_t graphExec1; + hipGraphNode_t hErrorNode_out; + 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(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); + 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); + + REQUIRE(hipErrorGraphExecUpdateFailure == ret); + REQUIRE(hipGraphExecUpdateErrorParametersChanged == updateResult_out); + REQUIRE(memcpy_B == hErrorNode_out); + HIP_CHECK(hipGraphDestroy(graph3)); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph1)); +} + +/** + * Test Description + * ------------------------ + * - Negative Scenario - 3 + Create graph1 and graph2 with different number node in it. + Instantiate graph1 in graphExec1 + Update graphExec1 with graph2 and verify. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_1") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; + hipError_t ret; + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec1; + hipGraphNode_t hErrorNode_out; + 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(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(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); + + REQUIRE(hipErrorGraphExecUpdateFailure == ret); +#if HT_NVIDIA + REQUIRE(hipGraphExecUpdateErrorNotSupported == updateResult_out); +#else + REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); +#endif + REQUIRE(NULL == hErrorNode_out); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); +} + +/** + * Test Description + * ------------------------ + * - Negative Scenario - + 4) Create a graph1 with 2 node and hipGraphInstantiate to create graphExec1 from it. + Delete a node from the Graph but not from its graphExec1 + Update graphExec1 with same graph (where node is deleted) and verify. + 5) Create a graph2 with 1 node and hipGraphInstantiate to create graphExec2 from it. + (Now graph1 and Graph2 have 1 node each with similar topology) + Update graphExec2 with graph1 (where node is deleted) and verify. + 6) Create a graph with 1 node & hipGraphInstantiate to create graphExec from it + Add one more node to the Graph Update graphExec with same graph + - (A node is deleted in hGraphExec but not its pair from hGraph) and verify + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Negative_Functional_CountDiffer_2") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; + hipError_t ret; + hipGraph_t graph1, graph2, graph3; + hipGraphExec_t graphExec1, graphExec2, graphExec3; + hipGraphNode_t hErrorNode_out; + 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(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); + REQUIRE(hipErrorGraphExecUpdateFailure == ret); + REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); +#if HT_NVIDIA + REQUIRE(NULL == hErrorNode_out); +#endif + } + 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(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + ret = hipGraphExecUpdate(graphExec2, graph1, &hErrorNode_out, + &updateResult_out); +#if HT_NVIDIA + REQUIRE(hipErrorGraphExecUpdateFailure == ret); + REQUIRE(hipGraphExecUpdateErrorNotSupported == updateResult_out); + REQUIRE(NULL == hErrorNode_out); +#else + REQUIRE(hipSuccess == ret); +#endif + HIP_CHECK(hipGraphDestroy(graph2)); + } + 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(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); + REQUIRE(hipErrorGraphExecUpdateFailure == ret); +#if HT_NVIDIA + REQUIRE(hipGraphExecUpdateErrorNotSupported == updateResult_out); +#else + REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); +#endif + REQUIRE(NULL == hErrorNode_out); + + HIP_CHECK(hipGraphDestroy(graph3)); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph1)); +} + +/** + * Test Description + * ------------------------ + * - Negative Scenario - + 7) Create a graph1 with memcpy_A, memcpy_B and memcpy_C, + add dependency as memcpy_A->memcpy_B->memcpy_C + and hipGraphInstantiate to create graphExec from it + Create a graph2 with same nodes and + dependency as memcpy_A->memcpy_C and memcpy_B->memcpy_C + and Update graphExec with graph2 and verify + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Negative_Dependent_NodesDiffer") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; + hipError_t ret; + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec; + hipGraphNode_t hErrorNode_out; + 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(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(hipGraphAddDependencies(graph2, &memcpy_A, &memcpy_C, 1)); + HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_B, &memcpy_C, 1)); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, + &updateResult_out); + + REQUIRE(hipErrorGraphExecUpdateFailure == ret); + REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); + REQUIRE(NULL != hErrorNode_out); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); +} + +/** + * Test Description + * ------------------------ + * - Negative Scenario - + 8) Create a graph1 with memcpy_A, memcpy_B and dependency memcpy_A->memcpy_B + and hipGraphInstantiate to create graphExec from it + Create a graph2 with memcpy_A, memsetNode and dependency memcpy_A->memsetNode + and Update graphExec with graph2 and verify + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Negative_NodeType_Changed") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphNode_t memcpy_A, memcpy_B, memsetNode; + hipError_t ret; + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec; + hipGraphNode_t hErrorNode_out; + 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(hipGraphAddDependencies(graph1, &memcpy_A, &memcpy_B, 1)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphCreate(&graph2, 0)); + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(C_d); + memsetParams.value = 3; + memsetParams.pitch = 0; + 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(hipGraphAddDependencies(graph2, &memcpy_A, &memsetNode, 1)); + ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out, + &updateResult_out); + REQUIRE(hipErrorGraphExecUpdateFailure == ret); +#if HT_NVIDIA + REQUIRE(hipGraphExecUpdateErrorTopologyChanged == updateResult_out); +#else + REQUIRE(hipGraphExecUpdateErrorNodeTypeChanged == updateResult_out); +#endif + REQUIRE(memsetNode == hErrorNode_out); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); +} + +/** + * Test Description + * ------------------------ + * - Negative Scenario - + 9) Multidevice case - set device 0 and + Create a graph1 with ketnelNode as vector_ADD + and hipGraphInstantiate to create graphExec from it + set device 1 and Create a graph2 with ketnelNode as vector_SUB + and Update graphExec with graph2 and verify. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Negative_MultiDevice_Context_Changed") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + size_t NElem{N}; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; + hipGraphNode_t kernel_vecADD, kernel_vecSUB; + hipError_t ret; + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec; + hipGraphNode_t hErrorNode_out; + hipGraphExecUpdateResult updateResult_out; + + int numDevices{}, peerAccess{}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + } + if (!peerAccess) { + WARN("Skipping test as peer device access is not found!"); + return; + } + HIP_CHECK(hipSetDevice(0)); + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + 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)); + hipKernelNodeParams kernelNodeParams{}; + 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(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)); + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); + + 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)); + memset(&kernelNodeParams, 0x00, sizeof(hipKernelNodeParams)); + 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(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); + + REQUIRE(hipErrorGraphExecUpdateFailure == ret); + REQUIRE(hipGraphExecUpdateErrorUnsupportedFunctionChange == updateResult_out); + REQUIRE(nullptr != hErrorNode_out); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * Test Description + * ------------------------ + * - Functional Scenario - + 1) Create a graph1 with ketnelNode as vector_ADD + and hipGraphInstantiate to create graphExec from it + Create a graph2 with ketnelNode as vector_SUB + and Update graphExec with graph2 and verify update should work as expected. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecUpdate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipGraphExecUpdate_Functional_KernelFunction_Changed") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + size_t NElem{N}; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; + hipGraphNode_t kernel_vecADD, kernel_vecSUB; + hipError_t ret; + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec; + hipGraphNode_t hErrorNode_out; + hipGraphExecUpdateResult updateResult_out; + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + 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)); + hipKernelNodeParams kernelNodeParams{}; + 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(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)); + memset(&kernelNodeParams, 0x00, sizeof(hipKernelNodeParams)); + 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(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); + REQUIRE(hipSuccess == ret); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Verify graph execution result + HipTest::checkVectorSUB(A_h, B_h, C_h, N); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipStreamDestroy(stream)); +}