From e9177bd83bd98aeaf0fede2d5c425cba6413cbdd Mon Sep 17 00:00:00 2001 From: Anavena Venkatesh Date: Mon, 9 Dec 2024 13:14:37 +0530 Subject: [PATCH] SWDEV-491363 - [catch2][dtest] Added test cases for missed graph APIs Change-Id: Ic740c52992b8d0a325f09b1db0adf599561e1a1d --- catch/hipTestMain/config/config_amd_linux | 6 +- catch/unit/errorHandling/CMakeLists.txt | 6 + .../unit/errorHandling/hipExtGetLastError.cc | 862 ++++++++++++++++++ catch/unit/graph/CMakeLists.txt | 8 +- catch/unit/graph/hipDrvGraphAddMemFreeNode.cc | 148 +++ .../hipDrvGraphExecMemcpyNodeSetParams.cc | 212 +++++ .../graph/hipDrvGraphMemcpyNodeGetParams.cc | 129 ++- catch/unit/graph/hipGraphExecGetFlags.cc | 149 +++ ...hipGraphExecMemcpyNodeSetParamsToSymbol.cc | 121 +-- catch/unit/graph/hipGraphExecNodeSetParams.cc | 156 ++++ catch/unit/graph/hipGraphNodeSetParams.cc | 163 ++++ 11 files changed, 1878 insertions(+), 82 deletions(-) create mode 100644 catch/unit/errorHandling/hipExtGetLastError.cc create mode 100644 catch/unit/graph/hipDrvGraphAddMemFreeNode.cc create mode 100644 catch/unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphExecGetFlags.cc create mode 100644 catch/unit/graph/hipGraphExecNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphNodeSetParams.cc diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index b07ee4e8b1..db8a99d2e4 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -39,8 +39,6 @@ "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters", "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", - "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic", - "Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Positive_Basic", "Unit_hipKernelNameRef_Negative_Parameters", "Unit_hipMemAdvise_No_Flag_Interference", "NOTE: The following 2 tests are disabled due to defect - EXSWHTEC-238", @@ -707,9 +705,6 @@ "Unit_hipFreeAsync_Negative_Parameters", "Unit_hipIpcGetMemHandle_Positive_Unique_Handles_Reused_Memory", "SWDEV-445928: These tests fail in PSDB stress test on 09/02/2024", - "Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint8_t", - "Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint16_t", - "Unit_hipGraphAddNodeTypeMemset_Positive_Basic - uint32_t", "Unit_hipCreateSurfaceObject_Negative_Parameters", "Unit_hipDestroySurfaceObject_Negative_Parameters", "Unit_Device___float2half_rd_Accuracy_Limited_Positive", @@ -757,6 +752,7 @@ "=== SWDEV-432554:Below test failed in stress test on 10/11/23 ===", "Unit_hipMemcpy3DAsync_Positive_Basic", "Unit_hipDrvMemcpy3DAsync_Positive_Basic", + "Print_Out_Attributes", "Unit_hipExtGetLinkTypeAndHopCount_Positive_Basic", "Unit_hipClock64_Positive_Basic", "Unit_hipClock_Positive_Basic", diff --git a/catch/unit/errorHandling/CMakeLists.txt b/catch/unit/errorHandling/CMakeLists.txt index 660740e2ad..09f543e07f 100644 --- a/catch/unit/errorHandling/CMakeLists.txt +++ b/catch/unit/errorHandling/CMakeLists.txt @@ -10,6 +10,12 @@ set(TEST_SRC ) add_executable(hipGetLastErrorEnv_Exe EXCLUDE_FROM_ALL hipGetLastErrorEnv_Exe.cc) add_executable(hipPeekAtLastErrorEnv_Exe EXCLUDE_FROM_ALL hipPeekAtLastErrorEnv_Exe.cc) +if(HIP_PLATFORM MATCHES "amd") + set(AMD_SRC + hipExtGetLastError.cc + ) + set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) +endif() hip_add_exe_to_target(NAME ErrorHandlingTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests diff --git a/catch/unit/errorHandling/hipExtGetLastError.cc b/catch/unit/errorHandling/hipExtGetLastError.cc new file mode 100644 index 0000000000..9ff71eea9c --- /dev/null +++ b/catch/unit/errorHandling/hipExtGetLastError.cc @@ -0,0 +1,862 @@ +/* +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 +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 +#include +#include +#include //NOLINT +#include //NOLINT + +static constexpr size_t WIDTH = 1024; +static constexpr size_t HEIGHT = 1024; +static constexpr size_t N = 1024 * 1024; +static constexpr size_t Nbytes = N * sizeof(int); + +/** +* @addtogroup hipExtGetLastError hipExtGetLastError +* @{ +* @ingroup ErrorHandlingTest +* `hipError_t hipExtGetLastError ( void )` - +* Returns the last error from a runtime call. +*/ + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMalloc api invalid arg call. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_Positive_Basic") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with ThreadedZigZagTest api call. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_Positive_Threaded") { + class HipGetLastErrorThreadedTest : public + ThreadedZigZagTest { + public: + void TestPart2() { + REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); + } + void TestPart3() { + HIP_CHECK(hipExtGetLastError()); + } + void TestPart4() { + REQUIRE_THREAD(hipExtGetLastError() == hipErrorInvalidValue); + } + }; + + HipGetLastErrorThreadedTest test; + test.run(); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpyPeerAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpyPeerAsync") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + + int can_access_peer = 0; + const auto src_device = 0; + const auto dst_device = 1; + + HIP_CHECK(hipSetDevice(src_device)); + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device)); + if (can_access_peer) { + HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0)); + + int *A_d, *B_d; + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + REQUIRE(A_d != nullptr); + + HIP_CHECK(hipSetDevice(dst_device)); + HIP_CHECK(hipMalloc(&B_d, Nbytes)); + REQUIRE(B_d != nullptr); + + HIP_CHECK(hipSetDevice(src_device)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpyPeerAsync(B_d, dst_device, A_d, src_device, + Nbytes * 2, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + HIP_CHECK(hipDeviceDisablePeerAccess(dst_device)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipSetDevice(dst_device)); + HIP_CHECK(hipFree(B_d)); + } else { + INFO("Peer access cannot be enabled between devices " << src_device << + " and devices " << dst_device); + } +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpyDtoHAsync api invalid arg call +*  Verify hipExtGetLastError status with hipMemcpyDtoDAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpyDtoHAsync") { + int *A_d, *B_d, *A_h; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, nullptr, nullptr, N, false); + + SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpyDtoHAsync(A_h, (hipDeviceptr_t)A_d, + Nbytes * 2, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + } + SECTION("Verify with hipMemcpyDtoDAsync api invalid arg call") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpyDtoDAsync((hipDeviceptr_t)A_d, + (hipDeviceptr_t)B_d, Nbytes * 2, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + } + + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, nullptr, nullptr, false); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpyParam2DAsync api invalid arg +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ +TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpyParam2DAsync") { + CHECK_IMAGE_SUPPORT + + float* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, * A_d{nullptr}; + size_t pitch_A; + size_t width{WIDTH * sizeof(float)}; + constexpr auto memsetval{100}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Allocating and Initializing the data + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, HEIGHT)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, &C_h, width*HEIGHT, false); + HipTest::setDefaultData(WIDTH*HEIGHT, A_h, B_h, C_h); + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, WIDTH, HEIGHT)); + + // Device to Host + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeDevice; + desc.dstMemoryType = hipMemoryTypeHost; + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = WIDTH; + desc.Height = HEIGHT; + + HIP_CHECK(hipExtGetLastError()); + desc.WidthInBytes = pitch_A+1; + HIP_CHECK_ERROR(hipMemcpyParam2DAsync(&desc, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipDrvMemcpy3DAsync api invalid arg +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipDrvMemcpy3DAsync") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipDrvMemcpy3DAsync(nullptr, hipStreamPerThread), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpy3DAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpy3DAsync") { + CHECK_IMAGE_SUPPORT + + constexpr int width{10}, height{10}, depth{10}; + auto size = width * height * depth * sizeof(int); + hipArray_t devArray; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + int *hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + + // Initialize host buffer + HipTest::setDefaultData(width*height*depth, hData, nullptr, nullptr); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, hipChannelFormatKindSigned); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, + make_hipExtent(width, height, 2), hipArrayDefault)); + + hipMemcpy3DParms myparams; + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, depth); + myparams.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparams.dstArray = devArray; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpy3DAsync(&myparams, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(devArray)); + HIP_CHECK(hipStreamDestroy(stream)); + free(hData); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpy2DToArrayAsync api invalid arg +*  Verify hipExtGetLastError status with hipMemcpy2DFromArrayAsync api invalid arg +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpy2D_To_From_ArrayAsync") { + int *hData = reinterpret_cast(malloc(WIDTH)); + REQUIRE(hData != nullptr); + memset(hData, 0, WIDTH); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, WIDTH, + WIDTH, HEIGHT, hipMemcpyHostToDevice, stream), + hipErrorInvalidHandle); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidHandle); + HIP_CHECK(hipExtGetLastError()); + } + SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(hData, WIDTH, nullptr, 0, 0, + WIDTH, HEIGHT, hipMemcpyDeviceToHost, stream), + hipErrorInvalidHandle); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidHandle); + HIP_CHECK(hipExtGetLastError()); + } + + HIP_CHECK(hipStreamDestroy(stream)); + free(hData); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipStreamAttachMemAsync api invalid arg +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipStreamAttachMemAsync") { + void* d_memory{nullptr}; + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemPrefetchAsync(reinterpret_cast(d_memory), + 0, hipMemAttachHost, 0), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipWaitExternalSemaphoresAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipWaitExternalSemaphoresAsync") { + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(nullptr, + &wait_params, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipSignalExternalSemaphoresAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipSignalExternalSemaphoresAsync") { + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(nullptr, + &signal_params, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemPrefetchAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemPrefetchAsync") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemPrefetchAsync(nullptr, 1024, 0), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpy2DAsync api invalid arg call +*  Verify hipExtGetLastError status with hipMemset2DAsync api invalid arg call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemcpy2DAsync") { + CHECK_IMAGE_SUPPORT + + int* A_h{nullptr}, *A_d{nullptr}; + size_t pitch_A; + size_t width{WIDTH * sizeof(int)}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Allocating memory + A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, WIDTH)); + REQUIRE(A_d != nullptr); + + // Initialize the data + HipTest::setDefaultData(WIDTH*HEIGHT, A_h, nullptr, nullptr); + + SECTION("Verify with hipMemcpy2DAsync api invalid arg call") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemcpy2DAsync(A_h, WIDTH*2, A_d, pitch_A, + WIDTH*sizeof(int), WIDTH, hipMemcpyDeviceToHost, stream), + hipErrorInvalidPitchValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidPitchValue); + HIP_CHECK(hipExtGetLastError()); + } + SECTION("Verify with hipMemset2DAsync api invalid arg call") { + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemset2DAsync(A_d, pitch_A, 22, + WIDTH*sizeof(int), WIDTH*9, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + } + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemsetAsync api invalid arg call. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipMemsetAsync") { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + REQUIRE(A_d != nullptr); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipMemsetAsync(A_d, 0, Nbytes * 2, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpyAsync api invalid arg call. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_MemCpyAsync") { + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, N); + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // testing to check error manually + HIP_CHECK_ERROR(hipMemcpyAsync(C_h, C_d, Nbytes+N, hipMemcpyDeviceToHost, 0), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipMemcpyAsync api invalid arg call +* Check in other thread this error should not report by hipExtGetLastError() +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +// Inside thread, both hipExtGetLastError() api call should not return error +static void thread_wait_func(int sleep_time) { + HIP_CHECK(hipExtGetLastError()); + std::this_thread::sleep_for(std::chrono::milliseconds(sleep_time * 1000)); + HIP_CHECK(hipExtGetLastError()); +} + +TEST_CASE("Unit_hipExtGetLastError_with_MemCpyAsync_thread") { + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, N); + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipExtGetLastError()); + + std::thread t(thread_wait_func, 2); + + // testing to check error manually + HIP_CHECK_ERROR(hipMemcpyAsync(C_h, C_d, Nbytes+N, hipMemcpyDeviceToHost, 0), + hipErrorInvalidValue); + + t.join(); + + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipGraphAddMemcpyNode1D api +* Create graph with one node as error consciously so it produces an error, +* which will be used to verify the behavior of hipExtGetLastError api. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipGraphAddMemcpyNode1D") { + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + size_t NElem{N}; + + hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, memcpy_E, kVecAdd; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + + 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(&stream)); + 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)); + + hipKernelNodeParams kNodeParams{}; + 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(&kVecAdd, graph, nullptr, 0, &kNodeParams)); + + // hipGraphAddMemcpyNode1D is called conciously with double the size + // so that it produces an error, which will be used to verify the + // behavior of hipExtGetLastError api. + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipGraphAddMemcpyNode1D(&memcpy_E, graph, nullptr, 0, C_h, + C_d, Nbytes * 2, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kVecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kVecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kVecAdd, &memcpy_C, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipExtGetLastError()); + 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(hipStreamDestroy(stream)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipStreamEndCapture api invalid arg call +* Create a graph1 with stream with ketnelNode as vector_ADD and +* hipStreamEndCapture on graph1 with hipGraphInstantiate to create graphExec +* Again hipStreamEndCapture on graph2 which will return hipErrorIllegalState +* now verify the behavior of hipExtGetLastError api with this call. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_with_hipStreamBegin_EndCapture") { + 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); + + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream)); + HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, &graph1)); + + // EndCapture is called conciously so that it produces an error, + // which will be used to verify the behavior of hipExtGetLastError api. + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph2), hipErrorIllegalState); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorIllegalState); + HIP_CHECK(hipExtGetLastError()); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); + 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(graph1)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with hipGraphCreate api invalid arg call. +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_error_check_with_hipGraphCreate") { + hipGraph_t graph; + hipError_t ret; + + HIP_CHECK(hipExtGetLastError()); + ret = hipGraphCreate(&graph, 1); + REQUIRE(ret == hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status should update with new api invalid arg call. +* Api hipGraphCreate -> return error hipErrorInvalidValue +* Api hipDeviceGetGraphMemAttribute -> return error hipErrorInvalidDevice +* Now hipExtGetLastError() api shoud return hipErrorInvalidDevice +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +TEST_CASE("Unit_hipExtGetLastError_success_before_error_check_again") { + int value = 0; + hipGraph_t graph; + + HIP_CHECK(hipExtGetLastError()); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidValue); + + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(-1, + hipGraphMemAttrUsedMemCurrent, &value), hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipExtGetLastError(), hipErrorInvalidDevice); + HIP_CHECK(hipExtGetLastError()); +} + +/** +* Test Description +* ------------------------ +*  - Verify hipExtGetLastError status with divide_by_zero kernel call +* Test source +* ------------------------ +*  - unit/errorHandling/hipExtGetLastError.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.4 +*/ + +static void __global__ devideKernl(int *i, int x, int y) { + *i = x/(x-y); +} + +TEST_CASE("Unit_hipExtGetLastError_with_Kernel_divide_by_zero") { + int *i_d; + int i = 9; + HIP_CHECK(hipMalloc(&i_d, sizeof(int))); + REQUIRE(i_d != nullptr); + HIP_CHECK(hipMemcpy(i_d, &i, sizeof(int), hipMemcpyHostToDevice)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipExtGetLastError()); + devideKernl<<<1, 1, 0, stream>>>(i_d, 3, 3); + HIP_CHECK(hipExtGetLastError()); + + HIP_CHECK(hipFree(i_d)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group hipExtGetLastError. + * @} + */ diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index 0b5cc8ecae..4376ea021f 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -85,13 +85,19 @@ set(TEST_SRC hipGraphAsyncUserObj.cc hipGraphExecBatchMemOpNodeSetParams.cc hipGraphAddBatchMemOpNode.cc - hipGraphBatchMemOpNodeSetParams.cc) + hipGraphBatchMemOpNodeSetParams.cc + hipDrvGraphExecMemcpyNodeSetParams.cc + hipDrvGraphAddMemFreeNode.cc) if(HIP_PLATFORM MATCHES "amd") set(AMD_SRC hipStreamCaptureExtModuleLaunchKernel.cc hipStreamBeginCaptureToGraph.cc hipGetProcAddressGraphApis.cc + # Below files are disbled in NVIDIA as PSDB builds are failing due to lower CUDA version. + hipGraphExecNodeSetParams.cc + hipGraphNodeSetParams.cc + hipGraphExecGetFlags.cc ) set(TEST_SRC ${TEST_SRC} ${AMD_SRC}) endif() diff --git a/catch/unit/graph/hipDrvGraphAddMemFreeNode.cc b/catch/unit/graph/hipDrvGraphAddMemFreeNode.cc new file mode 100644 index 0000000000..f9592a19f1 --- /dev/null +++ b/catch/unit/graph/hipDrvGraphAddMemFreeNode.cc @@ -0,0 +1,148 @@ +/* +Copyright (c) 2024 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 + +/** + * @addtogroup hipDrvGraphAddMemFreeNode hipDrvGraphAddMemFreeNode + * @{ + * @ingroup GraphTest + * `hipDrvGraphAddMemFreeNode ( hipGraphNode_t * phGraphNode, + * hipGraph_t hGraph, const hipGraphNode_t * dependencies, + * size_t numDependencies, hipDeviceptr_t dptr)'- + * Creates a memory free node and adds it to a graph. + */ +/** + * Test Description + * ------------------------ + * - Test to verify hipDrvGraphAddMemFreeNode behavior with invalid arguments: + * -# Null graph node + * -# Null graph + * -# Invalid numDependencies for null list of dependencies + * -# Invalid numDependencies and valid list for dependencies + * -# Null dev_ptr + * Test source + * ------------------------ + * - /unit/graph/hipDrvGraphAddMemFreeNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipDrvGraphAddMemFreeNode_Negative_Params") { + constexpr size_t N = 1024; + hipGraph_t graph; + hipGraphNode_t alloc_node, free_node; + std::vector dependencies; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipMemAllocNodeParams alloc_param; + memset(&alloc_param, 0, sizeof(alloc_param)); + alloc_param.bytesize = N; + alloc_param.poolProps.allocType = hipMemAllocationTypePinned; + alloc_param.poolProps.location.id = 0; + alloc_param.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&alloc_node, graph, nullptr, + 0, &alloc_param)); + REQUIRE(alloc_param.dptr != nullptr); + + SECTION("Passing nullptr to graph node") { + HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(nullptr, graph, + &alloc_node, 1, (hipDeviceptr_t)alloc_param.dptr), + hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph") { + HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&free_node, nullptr, + &alloc_node, 1, (hipDeviceptr_t)alloc_param.dptr), + hipErrorInvalidValue); + } + + SECTION("Pass invalid numDependencies") { + HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&free_node, graph, nullptr, + 5, (hipDeviceptr_t)alloc_param.dptr), + hipErrorInvalidValue); + } + + SECTION("Pass invalid numDependencies and valid list for dependencies") { + dependencies.push_back(alloc_node); + HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&free_node, graph, + dependencies.data(), dependencies.size() + 1, + (hipDeviceptr_t)alloc_param.dptr), + hipErrorInvalidValue); + } + + SECTION("Passing nullptr to dev_ptr") { + HIP_CHECK_ERROR(hipDrvGraphAddMemFreeNode(&alloc_node, graph, + &alloc_node, 1, 0), hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphDestroy(graph)); +} +/** + * Test Description + * ------------------------ + * - It will create memory alloation node and add to the graph then it + * will create memory free node to free allocated memory and add to the graph. + * Test source + * ------------------------ + * - /unit/graph/hipDrvGraphAddMemFreeNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 +*/ +TEST_CASE("Unit_hipDrvGraphAddMemFreeNode_Positive") { + constexpr size_t N = 1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipCtx_t context; + hipStream_t streamForGraph; + int deviceid = 0; + hipGraphNode_t node = nullptr, memFreeNode = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipSetDevice(deviceid)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipCtxCreate(&context, 0, deviceid)); + + hipMemAllocNodeParams alloc_param; + memset(&alloc_param, 0, sizeof(alloc_param)); + alloc_param.bytesize = N; + alloc_param.poolProps.allocType = hipMemAllocationTypePinned; + alloc_param.poolProps.location.id = 0; + alloc_param.poolProps.location.type = hipMemLocationTypeDevice; + HIP_CHECK(hipGraphAddMemAllocNode(&node, graph, nullptr, 0, &alloc_param)); + REQUIRE(alloc_param.dptr != nullptr); + + HIP_CHECK(hipDrvGraphAddMemFreeNode(&memFreeNode, graph, &node, 1, + (hipDeviceptr_t)alloc_param.dptr)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipCtxDestroy(context)); +} +/** +* End doxygen group GraphTest. +* @} +*/ diff --git a/catch/unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc b/catch/unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc new file mode 100644 index 0000000000..11445814e0 --- /dev/null +++ b/catch/unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc @@ -0,0 +1,212 @@ +/* +Copyright (c) 2024 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 + +/** + * @addtogroup hipDrvGraphExecMemcpyNodeSetParams hipDrvGraphExecMemcpyNodeSetParams + * @{ + * @ingroup GraphTest + * `hipDrvGraphExecMemcpyNodeSetParams((hipGraphExec_t hGraphExec, + * hipGraphNode_t hNode, const HIP_MEMCPY3D *copyParams, hipCt_t ctx)` - + * Sets the parameters for a memcpy node in the given graphExec + */ +/** + * Test Description + * ------------------------ + * - Verify API behavior with invalid arguments: + * -# graphExec is nullptr + * -# node is nullptr + * + * Test source + * ------------------------ + * - /unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipDrvGraphExecMemcpyNodeSetParams_Negative") { + size_t size = 10; + size_t numW = size * sizeof(int); + // Host Vectors + std::vector A_h(numW); + std::vector B_h(numW); + hipCtx_t context; + // Initialization + std::iota(A_h.begin(), A_h.end(), 1); + std::fill_n(B_h.begin(), size, 2); + int deviceid = 0; + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipSetDevice(deviceid)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipCtxCreate(&context, 0, deviceid)); + + + HIP_MEMCPY3D memCpy_params{}; + + memset(&memCpy_params, 0x0, sizeof(HIP_MEMCPY3D)); + memCpy_params.srcXInBytes = 0; + memCpy_params.srcY = 0; + memCpy_params.srcZ = 0; + memCpy_params.dstXInBytes = 0; + memCpy_params.dstY = 0; + memCpy_params.dstZ = 0; + memCpy_params.WidthInBytes = numW; + memCpy_params.Height = 1; + memCpy_params.Depth = 1; + memCpy_params.srcMemoryType = hipMemoryTypeHost; + memCpy_params.dstMemoryType = hipMemoryTypeHost; + memCpy_params.srcHost = A_h.data(); + memCpy_params.srcPitch = numW; + memCpy_params.srcHeight = 1; + memCpy_params.dstHost = B_h.data(); + memCpy_params.dstPitch = numW; + memCpy_params.dstHeight = 1; + + HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, + 0, &memCpy_params, context)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + SECTION("graphExec is nullptr") { + HIP_CHECK_ERROR(hipDrvGraphExecMemcpyNodeSetParams(nullptr, node, + &memCpy_params, context), hipErrorInvalidValue); + } + SECTION("node is nullptr") { + HIP_CHECK_ERROR(hipDrvGraphExecMemcpyNodeSetParams(graphExec, nullptr, + &memCpy_params, context), hipErrorInvalidValue); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipCtxDestroy(context)); +} +/** + * Test Description + * ------------------------ + * - It will verify soure data is copied destionation + * after adding new memcpy node. + * First will create and add mem copy node to graph. + * Create another mem copy param node with new source data and + * add to the graphExec using hipDrvGraphExecMemcpyNodeSetParams API and + * lauch the graph. Compare soure and destination data. + * + * Test source + * ------------------------ + * - /unit/graph/hipDrvGraphExecMemcpyNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ + +TEST_CASE("Unit_hipDrvGraphExecMemcpyNodeSetParams_Positive") { + size_t size = 10; + size_t numW = size * sizeof(int); + // Host Vectors + std::vector A_h(numW); + std::vector B_h(numW); + hipCtx_t context; + // Initialization + std::iota(A_h.begin(), A_h.end(), 1); + std::fill_n(B_h.begin(), size, 0); + int deviceid = 0; + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipSetDevice(deviceid)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipCtxCreate(&context, 0, deviceid)); + + + HIP_MEMCPY3D memCpy_params{}; + memset(&memCpy_params, 0x0, sizeof(HIP_MEMCPY3D)); + memCpy_params.srcXInBytes = 0; + memCpy_params.srcY = 0; + memCpy_params.srcZ = 0; + memCpy_params.dstXInBytes = 0; + memCpy_params.dstY = 0; + memCpy_params.dstZ = 0; + memCpy_params.WidthInBytes = numW; + memCpy_params.Height = 1; + memCpy_params.Depth = 1; + memCpy_params.srcMemoryType = hipMemoryTypeHost; + memCpy_params.dstMemoryType = hipMemoryTypeHost; + memCpy_params.srcHost = A_h.data(); + memCpy_params.srcPitch = numW; + memCpy_params.srcHeight = 1; + memCpy_params.dstHost = B_h.data(); + memCpy_params.dstPitch = numW; + memCpy_params.dstHeight = 1; + + HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, + 0, &memCpy_params, context)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + REQUIRE(memcmp(A_h.data(), B_h.data(), numW) == 0); + // Host Vectors for source + std::vector C_h(numW); + std::iota(C_h.begin(), C_h.end(), 10); + + HIP_MEMCPY3D memCpy_params2{}; + memset(&memCpy_params2, 0x0, sizeof(HIP_MEMCPY3D)); + memCpy_params2.srcXInBytes = 0; + memCpy_params2.srcY = 0; + memCpy_params2.srcZ = 0; + memCpy_params2.dstXInBytes = 0; + memCpy_params2.dstY = 0; + memCpy_params2.dstZ = 0; + memCpy_params2.WidthInBytes = numW; + memCpy_params2.Height = 1; + memCpy_params2.Depth = 1; + memCpy_params2.srcMemoryType = hipMemoryTypeHost; + memCpy_params2.dstMemoryType = hipMemoryTypeHost; + memCpy_params2.srcHost = C_h.data(); + memCpy_params2.srcPitch = numW; + memCpy_params2.srcHeight = 1; + memCpy_params2.dstHost = B_h.data(); + memCpy_params2.dstPitch = numW; + memCpy_params2.dstHeight = 1; + + HIP_CHECK(hipDrvGraphExecMemcpyNodeSetParams(graphExec, node, + &memCpy_params2, context)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + REQUIRE(memcmp(C_h.data(), B_h.data(), numW) == 0); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipCtxDestroy(context)); +} +/** +* End doxygen group GraphTest. +* @} +*/ + diff --git a/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc b/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc index 77595b11bb..d6e1c0e0a5 100644 --- a/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc +++ b/catch/unit/graph/hipDrvGraphMemcpyNodeGetParams.cc @@ -20,18 +20,17 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include #include +#include #include +#include -// hipDrvGraphAddMemcpyNode API is yet to be implemented in HIP runtime. -#if 0 /** * @addtogroup hipDrvGraphMemcpyNodeGetParams hipDrvGraphMemcpyNodeGetParams * @{ * @ingroup GraphTest - * `hipDrvGraphMemcpyNodeGetParams(hipGraphNode_t hNode, HIP_MEMCPY3D* nodeParams)` - - * Gets a memcpy node's parameters + * `hipDrvGraphMemcpyNodeGetParams(hipGraphNode_t hNode, HIP_MEMCPY3D* + * nodeParams)` - Gets a memcpy node's parameters * ________________________ * Test cases from other APIs: * - @ref Unit_hipDrvGraphMemcpyNodeSetParams_Positive_Basic @@ -63,37 +62,129 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeGetParams_Negative_Parameters") { LinearAllocGuard3D src_alloc(extent); LinearAllocGuard3D dst_alloc(extent); - auto params = - GetDrvMemcpy3DParms(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(), - make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToDevice); + auto params = GetDrvMemcpy3DParms( + dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToDevice); hipGraph_t graph = nullptr; hipGraphNode_t node = nullptr; SECTION("node == nullptr") { - HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(nullptr, ¶ms), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(nullptr, ¶ms), + hipErrorInvalidValue); } SECTION("pNodeParams == nullptr") { HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context)); - HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, nullptr), hipErrorInvalidValue); + HIP_CHECK( + hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context)); + HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, nullptr), + hipErrorInvalidValue); HIP_CHECK(hipGraphDestroy(graph)); } - +#if HT_AMD SECTION("Node is destroyed") { HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context)); + HIP_CHECK( + hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, ¶ms, context)); HIP_CHECK(hipGraphDestroy(graph)); - HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, ¶ms), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeGetParams(node, ¶ms), + hipErrorInvalidValue); } +#endif HIP_CHECK(hipCtxPopCurrent(&context)); HIP_CHECK(hipCtxDestroy(context)); } -#endif // if 0 - /** -* End doxygen group GraphTest. -* @} -*/ + * Test Description + * ------------------------ + * - Create graph node with memcopy parameters and add to graph, + * get the parameters from the created graph node and compare + * the all the values. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphMemcpyNodeGetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ + +TEST_CASE("Unit_hipDrvGraphMemcpyNodeGetParams_Positive") { + size_t size = 10; + size_t numW = size * sizeof(int); + // Host Vectors + std::vector A_h(numW); + std::vector B_h(numW); + hipCtx_t context; + // Initialization + std::iota(A_h.begin(), A_h.end(), 1); + std::fill_n(B_h.begin(), size, 0); + int deviceid = 0; + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipSetDevice(deviceid)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipCtxCreate(&context, 0, deviceid)); + + HIP_MEMCPY3D memCpy_params{}; + memset(&memCpy_params, 0x0, sizeof(HIP_MEMCPY3D)); + memCpy_params.srcXInBytes = 0; + memCpy_params.srcY = 0; + memCpy_params.srcZ = 0; + memCpy_params.dstXInBytes = 0; + memCpy_params.dstY = 0; + memCpy_params.dstZ = 0; + memCpy_params.WidthInBytes = numW; + memCpy_params.Height = 1; + memCpy_params.Depth = 1; + memCpy_params.srcMemoryType = hipMemoryTypeHost; + memCpy_params.dstMemoryType = hipMemoryTypeHost; + memCpy_params.srcHost = A_h.data(); + memCpy_params.srcPitch = numW; + memCpy_params.srcHeight = 1; + memCpy_params.dstHost = B_h.data(); + memCpy_params.dstPitch = numW; + memCpy_params.dstHeight = 1; + + HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, graph, nullptr, 0, &memCpy_params, + context)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + REQUIRE(memcmp(A_h.data(), B_h.data(), numW) == 0); + + HIP_MEMCPY3D memCpyGetParams{}; + HIP_CHECK(hipDrvGraphMemcpyNodeGetParams(node, &memCpyGetParams)); + + REQUIRE(memCpy_params.srcXInBytes == memCpyGetParams.srcXInBytes); + REQUIRE(memCpy_params.srcY == memCpyGetParams.srcY); + REQUIRE(memCpy_params.srcZ == memCpyGetParams.srcZ); + REQUIRE(memCpy_params.dstXInBytes == memCpyGetParams.dstXInBytes); + REQUIRE(memCpy_params.dstY == memCpyGetParams.dstY); + REQUIRE(memCpy_params.dstZ == memCpyGetParams.dstZ); + REQUIRE(memCpy_params.WidthInBytes == memCpyGetParams.WidthInBytes); + REQUIRE(memCpy_params.Height == memCpyGetParams.Height); + REQUIRE(memCpy_params.Depth == memCpyGetParams.Depth); + REQUIRE(memCpy_params.srcMemoryType == memCpyGetParams.srcMemoryType); + REQUIRE(memCpy_params.dstMemoryType == memCpyGetParams.dstMemoryType); + REQUIRE(memCpy_params.srcHost == memCpyGetParams.srcHost); + REQUIRE(memCpy_params.srcPitch == memCpyGetParams.srcPitch); + REQUIRE(memCpy_params.srcHeight == memCpyGetParams.srcHeight); + REQUIRE(memCpy_params.dstHost == memCpyGetParams.dstHost); + REQUIRE(memCpy_params.dstPitch == memCpyGetParams.dstPitch); + REQUIRE(memCpy_params.dstHeight == memCpyGetParams.dstHeight); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} +/** + * End doxygen group GraphTest. + * @} + */ diff --git a/catch/unit/graph/hipGraphExecGetFlags.cc b/catch/unit/graph/hipGraphExecGetFlags.cc new file mode 100644 index 0000000000..aeecaa3fe3 --- /dev/null +++ b/catch/unit/graph/hipGraphExecGetFlags.cc @@ -0,0 +1,149 @@ +/* +Copyright (c) 2024 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 +/** + * @addtogroup hipGraphExecGetFlags hipGraphExecGetFlags + * @{ + * @ingroup GraphTest + * `hipGraphExecGetFlags(hipGraphExec_t graphExec, + * unsigned long long *flags)` - + * Return the flags on executable graph + */ +/** + * Test Description + * ------------------------ + * - Verify API behavior with invalid arguments: + * -# graphExec is nullptr + * Test source + * ------------------------ + * - unit/graph/hipGraphExecGetFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphExecGetFlags_Negative") { + hipGraphExec_t graphExec; + unsigned long long flags; // NOLINT + constexpr size_t Nbytes = 10 * sizeof(int); + + hipGraphNode_t allocNodeA; + hipMemAllocNodeParams allocParam; + + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, + 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, + hipGraphInstantiateFlagAutoFreeOnLaunch)); + HIP_CHECK_ERROR(hipGraphExecGetFlags(nullptr, &flags), + hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Test Description + * ------------------------ + * - This test will verify the flags what we set while initiating + * graph will be matching with flags values getting from + * the hipGraphExecGetFlags API call. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecGetFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphExecGetFlags_positive") { + hipGraphExec_t graphExec; + unsigned long long flags; // NOLINT + hipGraph_t graph; + constexpr size_t Nbytes = 10 * sizeof(int); + + hipGraphNode_t allocNodeA; + hipMemAllocNodeParams allocParam; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, + 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + + SECTION("flag is 0") { + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, + 0)); + + HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags)); + REQUIRE(flags == 0); + } + SECTION("flag is hipGraphInstantiateFlagAutoFreeOnLaunch") { + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, + hipGraphInstantiateFlagAutoFreeOnLaunch)); + + HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags)); + REQUIRE(flags == hipGraphInstantiateFlagAutoFreeOnLaunch); + } + +// The below feature flags are not implemented +// hipGraphInstantiateFlagUpload +// hipGraphInstantiateFlagDeviceLaunch +// hipGraphInstantiateFlagUseNodePriority +#if 0 + SECTION("flag is hipGraphInstantiateFlagUpload") { + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, + hipGraphInstantiateFlagUpload)); + + HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags)); + REQUIRE(flags == hipGraphInstantiateFlagUpload); + } + + SECTION("flag is hipGraphInstantiateFlagDeviceLaunch") { + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, + hipGraphInstantiateFlagDeviceLaunch)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + REQUIRE(flags == hipGraphInstantiateFlagDeviceLaunch); + } + SECTION("flag is hipGraphInstantiateFlagUseNodePriority") { + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, + hipGraphInstantiateFlagUseNodePriority)); + HIP_CHECK(hipGraphExecGetFlags(graphExec, &flags)); + REQUIRE(flags == hipGraphInstantiateFlagUseNodePriority); + } +#endif + HIP_CHECK(hipGraphDestroy(graph)); +} +/** +* End doxygen group GraphTest. +* @} +*/ + diff --git a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc index ee05f191b9..78fd4aa13b 100644 --- a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc +++ b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc @@ -19,12 +19,10 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include +#include #include #include - -#include -#include - #include "graph_memcpy_to_from_symbol_common.hh" HIP_GRAPH_MEMCPY_FROM_SYMBOL_NODE_DEFINE_GLOBALS(char) @@ -38,24 +36,27 @@ HIP_GRAPH_MEMCPY_FROM_SYMBOL_NODE_DEFINE_ALTERNATE_GLOBALS(float) HIP_GRAPH_MEMCPY_FROM_SYMBOL_NODE_DEFINE_ALTERNATE_GLOBALS(double) template -void GraphExecMemcpyToSymbolSetParamsShell(const void* symbol, const void* alt_symbol, - size_t offset, const std::vector set_values) { - const auto f = [alt_symbol, is_arr = set_values.size() > 1](const void* symbol, void* src, - size_t count, size_t offset, - hipMemcpyKind direction) { +void GraphExecMemcpyToSymbolSetParamsShell(const void *symbol, + const void *alt_symbol, + size_t offset, + const std::vector set_values) { + const auto f = [alt_symbol, is_arr = set_values.size() > 1]( + const void *symbol, void *src, size_t count, size_t offset, + hipMemcpyKind direction) { hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); hipGraphNode_t node = nullptr; HIP_CHECK(hipGraphAddMemcpyNodeToSymbol( - &node, graph, nullptr, 0, alt_symbol, reinterpret_cast(src) + is_arr, - count - is_arr * sizeof(T), offset + is_arr * sizeof(T), direction)); + &node, graph, nullptr, 0, alt_symbol, + reinterpret_cast(src) + is_arr, count - is_arr * sizeof(T), + offset + is_arr * sizeof(T), direction)); hipGraphExec_t graph_exec = nullptr; HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphExecMemcpyNodeSetParamsToSymbol(graph_exec, node, symbol, src, count, offset, - direction)); + HIP_CHECK(hipGraphExecMemcpyNodeSetParamsToSymbol( + graph_exec, node, symbol, src, count, offset, direction)); HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); @@ -70,26 +71,30 @@ void GraphExecMemcpyToSymbolSetParamsShell(const void* symbol, const void* alt_s } /** - * @addtogroup hipGraphExecMemcpyNodeSetParamsToSymbol hipGraphExecMemcpyNodeSetParamsToSymbol + * @addtogroup hipGraphExecMemcpyNodeSetParamsToSymbol + * hipGraphExecMemcpyNodeSetParamsToSymbol * @{ * @ingroup GraphTest - * `hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node, - * const void *symbol, void *src, size_t count, size_t offset, hipMemcpyKind kind)` - - * Sets the parameters for a memcpy node in the given graphExec to copy to a symbol on the device + * `hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, + * hipGraphNode_t node, const void *symbol, void *src, size_t count, size_t + * offset, hipMemcpyKind kind)` - Sets the parameters for a memcpy node in the + * given graphExec to copy to a symbol on the device */ /** * Test Description * ------------------------ - * - Verify that data is correctly copied to a symbol after node parameters are set following - * node addition. A graph is constructed to which a MemcpyToSymbol node is added with valid but - * incorrect parameters. After the graph is instantiated the parameters are updated to correct - * values and the graph executed. After graph execution, a MemcpyFromSymbol is performed and the - * copied values are compared against values known to have been copied to symbol memory previously. - * The test is run for scalar, const scalar, array, and const array symbols of types char, int, - * float and double. For array symbols, the test is repeated for zero and non-zero offset values. - * Verification is performed for destination memory allocated on host and device. - * Test source + * - Verify that data is correctly copied to a symbol after node parameters + * are set following node addition. A graph is constructed to which a + * MemcpyToSymbol node is added with valid but incorrect parameters. After the + * graph is instantiated the parameters are updated to correct values and the + * graph executed. After graph execution, a MemcpyFromSymbol is performed and + * the copied values are compared against values known to have been copied to + * symbol memory previously. The test is run for scalar, const scalar, array, + * and const array symbols of types char, int, float and double. For array + * symbols, the test is repeated for zero and non-zero offset values. + * Verification is performed for destination memory allocated on host and + * device. Test source * ------------------------ * - unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc * Test requirements @@ -98,23 +103,23 @@ void GraphExecMemcpyToSymbolSetParamsShell(const void* symbol, const void* alt_s */ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic") { SECTION("char") { - HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10, - char); + HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST( + GraphExecMemcpyToSymbolSetParamsShell, 10, char); } SECTION("int") { - HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10, - int); + HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST( + GraphExecMemcpyToSymbolSetParamsShell, 10, int); } SECTION("float") { - HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10, - float); + HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST( + GraphExecMemcpyToSymbolSetParamsShell, 10, float); } SECTION("double") { - HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST(GraphExecMemcpyToSymbolSetParamsShell, 10, - double); + HIP_GRAPH_MEMCPY_NODE_SET_PARAMS_TO_FROM_SYMBOL_TEST( + GraphExecMemcpyToSymbolSetParamsShell, 10, double); } } @@ -132,8 +137,8 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic") { * -# kind is illogical (hipMemcpyDeviceToHost) * -# kind is an invalid enum value * -# Changing memcpy direction - * -# Changing src to memory allocated on a different device than the original src - * Test source + * -# Changing src to memory allocated on a different device than the + * original src Test source * ------------------------ * - unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc * Test requirements @@ -147,35 +152,37 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") { LinearAllocGuard var(LinearAllocs::hipMalloc, sizeof(int)); hipGraphNode_t node = nullptr; - HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&node, graph, nullptr, 0, SYMBOL(int_device_var), - var.ptr(), sizeof(*var.ptr()), 0, hipMemcpyDefault)); + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol( + &node, graph, nullptr, 0, SYMBOL(int_device_var), var.ptr(), + sizeof(*var.ptr()), 0, hipMemcpyDefault)); hipGraphExec_t graph_exec = nullptr; HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); SECTION("hGraphExec == nullptr") { - HIP_CHECK_ERROR( - hipGraphExecMemcpyNodeSetParamsToSymbol(nullptr, node, SYMBOL(int_device_var), var.ptr(), - sizeof(*var.ptr()), 0, hipMemcpyDefault), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol( + nullptr, node, SYMBOL(int_device_var), var.ptr(), + sizeof(*var.ptr()), 0, hipMemcpyDefault), + hipErrorInvalidValue); } SECTION("node == nullptr") { - HIP_CHECK_ERROR( - hipGraphExecMemcpyNodeSetParamsToSymbol(graph_exec, nullptr, SYMBOL(int_device_var), - var.ptr(), sizeof(*var.ptr()), 0, hipMemcpyDefault), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol( + graph_exec, nullptr, SYMBOL(int_device_var), var.ptr(), + sizeof(*var.ptr()), 0, hipMemcpyDefault), + hipErrorInvalidValue); } MemcpyToSymbolCommonNegative( - std::bind(hipGraphExecMemcpyNodeSetParamsToSymbol, graph_exec, node, _1, _2, _3, _4, _5), + std::bind(hipGraphExecMemcpyNodeSetParamsToSymbol, graph_exec, node, _1, + _2, _3, _4, _5), SYMBOL(int_device_var), var.ptr(), sizeof(*var.ptr())); SECTION("Changing memcpy direction") { - HIP_CHECK_ERROR( - hipGraphExecMemcpyNodeSetParamsToSymbol(graph_exec, node, SYMBOL(int_device_var), var.ptr(), - sizeof(*var.ptr()), 0, hipMemcpyHostToDevice), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol( + graph_exec, node, SYMBOL(int_device_var), var.ptr(), + sizeof(*var.ptr()), 0, hipMemcpyHostToDevice), + hipErrorInvalidValue); } SECTION("Changing src allocation device") { @@ -185,9 +192,9 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") { HIP_CHECK(hipSetDevice(1)); LinearAllocGuard new_var(LinearAllocs::hipMalloc, sizeof(int)); HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParamsToSymbol( - graph_exec, node, SYMBOL(int_device_var), new_var.ptr(), - sizeof(*new_var.ptr()), 0, static_cast(-1)), - hipErrorInvalidValue); + graph_exec, node, SYMBOL(int_device_var), + new_var.ptr(), sizeof(*new_var.ptr()), + 0, hipMemcpyDefault), hipErrorInvalidValue); } } @@ -196,6 +203,6 @@ TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative_Parameters") { } /** -* End doxygen group GraphTest. -* @} -*/ + * End doxygen group GraphTest. + * @} + */ diff --git a/catch/unit/graph/hipGraphExecNodeSetParams.cc b/catch/unit/graph/hipGraphExecNodeSetParams.cc new file mode 100644 index 0000000000..bac926fe7c --- /dev/null +++ b/catch/unit/graph/hipGraphExecNodeSetParams.cc @@ -0,0 +1,156 @@ +/* +Copyright (c) 2024 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 + +/** + * @addtogroup hipGraphExecNodeSetParams hipGraphExecNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphExecNodeSetParams(hipGraphExec_t graphExec, hipGraphNode_t node, + * hipGraphNodeParams *nodeParams)` - + * Updates parameters of a created node on executable graph + */ + +/** + * Test Description + * ------------------------ + * - Verify API behavior with invalid arguments: + * -# gGraphExec is nullptr + * -# node is nullptr + * -# nodeParams is nullptr + * Test source + * ------------------------ + * - unit/graph/hipGraphExecNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphExecNodeSetParams_Negative_Parameters") { + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + char *A_d; + size_t Nbytes = 10 * sizeof(char); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + node_params.type = hipGraphNodeTypeMemset; + node_params.memset.dst = A_d; + node_params.memset.elementSize = sizeof(char); + node_params.memset.width = 10; + node_params.memset.height = 1; + node_params.memset.pitch = 10; + node_params.memset.value = 99; + + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + SECTION("hGraphExec == nullptr") { + HIP_CHECK_ERROR( + hipGraphExecNodeSetParams(nullptr, node, &node_params), + hipErrorInvalidValue); + } + + SECTION("node == nullptr") { + HIP_CHECK_ERROR( + hipGraphExecNodeSetParams(graphExec, nullptr, &node_params), + hipErrorInvalidValue); + } + + SECTION("node params == nullptr") { + HIP_CHECK_ERROR(hipGraphExecNodeSetParams(graphExec, node, nullptr), + hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - This will verify the new node param values are successfully + * copied to graph node, after launching graphExec + * Test source + * ------------------------ + * - unit/graph/hipGraphExecNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphExecNodeSetParams_Positive") { + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + char *A_d = nullptr, *A_h = nullptr; + size_t Nbytes = 10 * sizeof(char); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + node_params.type = hipGraphNodeTypeMemset; + node_params.memset.dst = A_d; + node_params.memset.elementSize = sizeof(char); + node_params.memset.width = 10; + node_params.memset.height = 1; + node_params.memset.pitch = 10; + node_params.memset.value = 99; + + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < 10; i++) { + REQUIRE(A_h[i] == 99); + } + + hipGraphNodeParams node_params2 = {}; + node_params2.type = hipGraphNodeTypeMemset; + node_params2.memset.dst = A_d; + node_params2.memset.elementSize = sizeof(char); + node_params2.memset.width = 10; + node_params2.memset.height = 1; + node_params2.memset.pitch = 10; + node_params2.memset.value = 110; + + HIP_CHECK(hipGraphExecNodeSetParams(graphExec, node, &node_params2)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < 10; i++) { + REQUIRE(A_h[i] == 110); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(A_d)); + free(A_h); +} +/** +* End doxygen group GraphTest. +* @} +*/ diff --git a/catch/unit/graph/hipGraphNodeSetParams.cc b/catch/unit/graph/hipGraphNodeSetParams.cc new file mode 100644 index 0000000000..745f352ae0 --- /dev/null +++ b/catch/unit/graph/hipGraphNodeSetParams.cc @@ -0,0 +1,163 @@ +/* +Copyright (c) 2024 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 + +/** + * @addtogroup hipGraphNodeSetParams hipGraphNodeSetParams + * @{ + * @ingroup GraphTest + * `hipGraphNodeSetParams(hipGraphNode_t node, + * hipGraphNodeParams *nodeParams)` - + * Updates parameters of a graph’s node + */ +/** + * Test Description + * ------------------------ + * - Verify API behavior with invalid arguments: + * -# node is nullptr + * -# nodeParams is nullptr + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphNodeSetParams_Negative_Parameters") { + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + char *A_d; + char *A_h; + size_t N = 10; + size_t Nbytes = N * sizeof(char); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + node_params.type = hipGraphNodeTypeMemset; + node_params.memset.dst = A_d; + node_params.memset.elementSize = sizeof(char); + node_params.memset.width = N; + node_params.memset.height = 1; + node_params.memset.pitch = N; + node_params.memset.value = 99; + + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + hipGraphNodeParams node_params2 = {}; + node_params2.type = hipGraphNodeTypeMemset; + node_params2.memset.dst = A_d; + node_params2.memset.elementSize = sizeof(char); + node_params2.memset.width = N; + node_params2.memset.height = 1; + node_params2.memset.pitch = N; + node_params2.memset.value = 110; + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphNodeSetParams(nullptr, &node_params2), + hipErrorInvalidValue); + } + + SECTION("nodeParams == nullptr") { + HIP_CHECK_ERROR(hipGraphNodeSetParams(node, nullptr), + hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Test Description + * ------------------------ + * - This will verify the new node param values are successfully + * copied to graph node + * Test source + * ------------------------ + * - unit/graph/hipGraphNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphNodeSetParams_Positive") { + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t node; + hipGraphNodeParams node_params = {}; + char *A_d, *A_h = nullptr; + size_t N = 10; + size_t Nbytes = N * sizeof(char); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + node_params.type = hipGraphNodeTypeMemset; + node_params.memset.dst = A_d; + node_params.memset.elementSize = sizeof(char); + node_params.memset.width = N; + node_params.memset.height = 1; + node_params.memset.pitch = N; + node_params.memset.value = 99; + + HIP_CHECK(hipGraphAddNode(&node, graph, nullptr, 0, &node_params)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) { + REQUIRE(A_h[i] == 99); + } + + hipGraphNodeParams node_params2 = {}; + node_params2.type = hipGraphNodeTypeMemset; + node_params2.memset.dst = A_d; + node_params2.memset.elementSize = sizeof(char); + node_params2.memset.width = N; + node_params2.memset.height = 1; + node_params2.memset.pitch = N; + node_params2.memset.value = 110; + + HIP_CHECK(hipGraphNodeSetParams(node, &node_params2)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) { + REQUIRE(A_h[i] == 110); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(A_d)); + free(A_h); +} + +/** +* End doxygen group GraphTest. +* @} +*/