From 9a93994fcf961d719148c9635a10fca5e31a1fdf Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 16 Jan 2023 16:30:50 +0530 Subject: [PATCH 1/3] SWDEV-306122 - [catch2][dtest] Modify some negative test hipGraphMemcpyNodeSetParams1D (#76) Change-Id: Ie75c1b55b35efda98db77c425e73d495c7e0083d --- catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc b/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc index 92f51ad3aa..86e439e528 100644 --- a/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc +++ b/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc @@ -77,18 +77,18 @@ TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Negative") { #if HT_AMD SECTION("Pass same pointer as source ptr and destination ptr") { ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d, Nbytes, - hipMemcpyHostToDevice); + hipMemcpyDeviceToDevice); REQUIRE(hipErrorInvalidValue == ret); } +#endif SECTION("Pass overlap memory where destination ptr is ahead of source ptr") { ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d-5, Nbytes, - hipMemcpyHostToDevice); - REQUIRE(hipSuccess == ret); + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); } -#endif SECTION("Pass overlap memory where source ptr is ahead of destination ptr") { ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d+5, A_d, Nbytes-5, - hipMemcpyHostToDevice); + hipMemcpyDeviceToDevice); REQUIRE(hipErrorInvalidValue == ret); } SECTION("Copy more than allocated memory") { From b8a83b3cc476532c60078702613da903150f9e5f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 16 Jan 2023 16:31:04 +0530 Subject: [PATCH 2/3] SWDEV-306122 - [catch2][dtest] hipGraphCreate, hipGraphDestroy and hipGraphExecDestroy APIs test (#79) Change-Id: I8219122c6c0f85a70a22f66037c0f2920047ab8b --- catch/unit/graph/CMakeLists.txt | 3 + catch/unit/graph/hipGraphCreate.cc | 44 ++++++++++++++ catch/unit/graph/hipGraphDestroy.cc | 81 +++++++++++++++++++++++++ catch/unit/graph/hipGraphExecDestroy.cc | 73 ++++++++++++++++++++++ 4 files changed, 201 insertions(+) create mode 100644 catch/unit/graph/hipGraphCreate.cc create mode 100644 catch/unit/graph/hipGraphDestroy.cc create mode 100644 catch/unit/graph/hipGraphExecDestroy.cc diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index 2991263331..6974cb81ae 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -83,6 +83,9 @@ set(TEST_SRC hipUserObjectCreate.cc hipGraphDebugDotPrint.cc hipGraphCloneComplx.cc + hipGraphCreate.cc + hipGraphDestroy.cc + hipGraphExecDestroy.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/catch/unit/graph/hipGraphCreate.cc b/catch/unit/graph/hipGraphCreate.cc new file mode 100644 index 0000000000..befcc69c7e --- /dev/null +++ b/catch/unit/graph/hipGraphCreate.cc @@ -0,0 +1,44 @@ +/* +Copyright (c) 2022 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 + +/** +Negative Testcase Scenarios : +1) Creating HipGraph with nullptr. +2) Creating hipGraph with non-zero mode. +*/ + +TEST_CASE("Unit_hipGraphCreate_Negative") { + hipError_t ret; + SECTION("Creating HipGraph with nullptr") { + ret = hipGraphCreate(nullptr, 0); + REQUIRE(hipErrorInvalidValue == ret); + } + + SECTION("Creating hipGraph with non-zero mode") { + hipGraph_t graph{}; + ret = hipGraphCreate(&graph, -1); + REQUIRE(hipErrorInvalidValue == ret); + + ret = hipGraphCreate(nullptr, -1); + REQUIRE(hipErrorInvalidValue == ret); + } +} + diff --git a/catch/unit/graph/hipGraphDestroy.cc b/catch/unit/graph/hipGraphDestroy.cc new file mode 100644 index 0000000000..15b8263f13 --- /dev/null +++ b/catch/unit/graph/hipGraphDestroy.cc @@ -0,0 +1,81 @@ +/* +Copyright (c) 2022 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 + +/** +Negative Testcase Scenarios : +1) Pass graph as nullptr and verify. +2) Destroy already destroyed graph and check api returns error code. +3) Destroy graph when is in use and make sure api handles it gracefully. +*/ + +TEST_CASE("Unit_hipGraphDestroy_Negative") { + hipError_t ret; + SECTION("Deleting HipGraph with nullptr") { + ret = hipGraphDestroy(nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_AMD + SECTION("Destroy already destroyed graph and check api returns error code") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphDestroy(graph)); + + ret = hipGraphDestroy(graph); + REQUIRE(hipErrorIllegalState == ret); + } +#endif + SECTION("Destroy graph when is in use and make sure api handles" + " it gracefully.") { + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + char *devData; + HIP_CHECK(hipMalloc(&devData, 1024)); + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = 1024; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(graphExec != nullptr); + + HIP_CHECK(hipGraphDestroy(graph)); + + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + } +} + diff --git a/catch/unit/graph/hipGraphExecDestroy.cc b/catch/unit/graph/hipGraphExecDestroy.cc new file mode 100644 index 0000000000..7ed2e0d718 --- /dev/null +++ b/catch/unit/graph/hipGraphExecDestroy.cc @@ -0,0 +1,73 @@ +/* +Copyright (c) 2022 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 + +/** +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. +*/ + +TEST_CASE("Unit_hipGraphExecDestroy_Negative") { + hipError_t ret; + SECTION("Pass hipGraphExecDestroy with nullptr") { + ret = hipGraphExecDestroy(nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipGraphExecDestroy with un-initilze structure") { + hipGraphExec_t graphExec{}; + ret = hipGraphExecDestroy(graphExec); + REQUIRE(hipErrorInvalidValue == ret); + } +} + +TEST_CASE("Unit_hipGraphExecDestroy_Sequence") { + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + char *devData; + HIP_CHECK(hipMalloc(&devData, 1024)); + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = 1024; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(graphExec != nullptr); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + From fb919e7995cb7f6e3f2feab7f43c0b1bcdc16f64 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 16 Jan 2023 16:31:47 +0530 Subject: [PATCH 3/3] SWDEV-306122 - [catch2][dtest] Added test for hipGraphUpload API (#89) Change-Id: If0cd652926c1308d8185648516256e13a7b3e5fd --- catch/unit/graph/CMakeLists.txt | 1 + catch/unit/graph/hipGraphUpload.cc | 236 +++++++++++++++++++++++++++++ 2 files changed, 237 insertions(+) create mode 100644 catch/unit/graph/hipGraphUpload.cc diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index 6974cb81ae..c9d1bb34ea 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -86,6 +86,7 @@ set(TEST_SRC hipGraphCreate.cc hipGraphDestroy.cc hipGraphExecDestroy.cc + hipGraphUpload.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/catch/unit/graph/hipGraphUpload.cc b/catch/unit/graph/hipGraphUpload.cc new file mode 100644 index 0000000000..d17fb27236 --- /dev/null +++ b/catch/unit/graph/hipGraphUpload.cc @@ -0,0 +1,236 @@ +/* +Copyright (c) 2022 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 + +/** + * Functional Test for API - hipGraphUpload + - Make graph from hipStreamBeginCapture with a stream. + Upload the graph into different stream and execute the graph and verify. + */ + +static void hipGraphUploadFunctional_with_hipStreamBeginCapture( + hipStream_t iStream) { + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1024; + size_t Nbytes = N * sizeof(float); + + int *A_d, *C_d; + int *A_h, *C_h; + HipTest::initArrays(&A_d, nullptr, &C_d, &A_h, nullptr, &C_h, N, false); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + + // Validate end capture is successful + REQUIRE(graph != nullptr); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(graphExec != nullptr); + + HIP_CHECK(hipGraphUpload(graphExec, iStream)); + HIP_CHECK(hipGraphLaunch(graphExec, iStream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamSynchronize(iStream)); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + UNSCOPED_INFO("A and C not matching at " << i); + } + } + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); +} + +/** + * Functional Test for API - hipGraphUpload + - Make graph by creating graph and add node and functionality to it. + Upload the graph into different stream and execute the graph and verify. + */ + +static void hipGraphUploadFunctional_with_stream(hipStream_t stream) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, kNodeAdd; + hipKernelNodeParams kNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + + 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(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)); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kNodeAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kNodeAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kNodeAdd, &memcpy_C, 1)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + + HIP_CHECK(hipGraphUpload(graphExec, stream)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Verify graph execution result + HipTest::checkVectorADD(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(graph)); +} + +TEST_CASE("Unit_hipGraphUpload_Functional") { + SECTION("Pass a stream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + hipGraphUploadFunctional_with_hipStreamBeginCapture(stream); + hipGraphUploadFunctional_with_stream(stream); + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Pass stream as default stream") { + hipGraphUploadFunctional_with_hipStreamBeginCapture(0); + hipGraphUploadFunctional_with_stream(0); + } + SECTION("Pass stream as hipStreamPerThread") { + hipGraphUploadFunctional_with_hipStreamBeginCapture(hipStreamPerThread); + hipGraphUploadFunctional_with_stream(hipStreamPerThread); + } +} + +TEST_CASE("Unit_hipGraphUpload_Functional_multidevice_test") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + if (numDevices > 0) { + SECTION("Pass a common stream for all device") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipSetDevice(i)); + hipGraphUploadFunctional_with_hipStreamBeginCapture(stream); + hipGraphUploadFunctional_with_stream(stream); + } + + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Pass a separate stream for each device") { + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipSetDevice(i)); + + hipStream_t dStream; + HIP_CHECK(hipStreamCreate(&dStream)); + hipGraphUploadFunctional_with_hipStreamBeginCapture(dStream); + hipGraphUploadFunctional_with_stream(dStream); + HIP_CHECK(hipStreamDestroy(dStream)); + } + } + SECTION("Pass stream as default stream for each device") { + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipSetDevice(i)); + hipGraphUploadFunctional_with_hipStreamBeginCapture(0); + hipGraphUploadFunctional_with_stream(0); + } + } + SECTION("Pass stream as hipStreamPerThread for each device") { + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipSetDevice(i)); + hipGraphUploadFunctional_with_hipStreamBeginCapture(hipStreamPerThread); + hipGraphUploadFunctional_with_stream(hipStreamPerThread); + } + } + } else { + SUCCEED("Skipped the testcase as there is no device to test."); + } +} + +/** +* Negative Test for API - hipGraphUpload Argument Check +1) Pass graphExec node as nullptr. +2) Pass graphExec node as uninitialize object +3) Pass stream as uninitialize object +*/ + +TEST_CASE("Unit_hipGraphUpload_Negative_Argument_Check") { + hipGraphExec_t graphExec{}; + hipError_t ret; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Pass graphExec node as nullptr") { + ret = hipGraphUpload(nullptr, stream); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass graphExec node as uninitialize object") { + ret = hipGraphUpload(graphExec, stream); + REQUIRE(hipErrorInvalidValue == ret); + } + 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)); + + ret = hipGraphUpload(graphExec, stream1); + REQUIRE(hipSuccess == ret); + } + HIP_CHECK(hipStreamDestroy(stream)); +} +