From 9d62df85a16a0e89130597d57c81f34e79e76891 Mon Sep 17 00:00:00 2001 From: nives-vukovic <110852104+nives-vukovic@users.noreply.github.com> Date: Tue, 17 Jan 2023 12:56:17 +0100 Subject: [PATCH] EXSWHTEC-91 - Implement tests for memcpy of 1D/2D hipArray (#32) - Implement tests for hipMemcpyAtoH using resource guards and templates - Implement tests for hipMemcpyHtoA using resource guards and templates - Implement tests for hipMemcpy2DFromArray using resource guards and templates - Implement tests for hipMemcpy2DToArray using resource guards and templates --- catch/unit/memory/CMakeLists.txt | 8 + catch/unit/memory/hipMemcpy2DFromArray.cc | 515 ++++++++---------- catch/unit/memory/hipMemcpy2DFromArray_old.cc | 331 +++++++++++ catch/unit/memory/hipMemcpy2DToArray.cc | 511 ++++++++--------- catch/unit/memory/hipMemcpy2DToArray_old.cc | 332 +++++++++++ catch/unit/memory/hipMemcpyAtoH.cc | 263 +++------ catch/unit/memory/hipMemcpyAtoH_old.cc | 219 ++++++++ catch/unit/memory/hipMemcpyHtoA.cc | 268 +++------ catch/unit/memory/hipMemcpyHtoA_old.cc | 229 ++++++++ 9 files changed, 1720 insertions(+), 956 deletions(-) create mode 100644 catch/unit/memory/hipMemcpy2DFromArray_old.cc create mode 100644 catch/unit/memory/hipMemcpy2DToArray_old.cc create mode 100644 catch/unit/memory/hipMemcpyAtoH_old.cc create mode 100644 catch/unit/memory/hipMemcpyHtoA_old.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index d0dba71e6c..c89fae9bfc 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -25,6 +25,7 @@ set(TEST_SRC memset.cc malloc.cc hipMemcpy2DToArray.cc + hipMemcpy2DToArray_old.cc hipMemcpy2DToArrayAsync.cc hipMemcpy2DToArrayAsync_old.cc hipMemcpy3D.cc @@ -34,10 +35,13 @@ set(TEST_SRC hipMemcpy2D.cc hipMemcpy2DAsync.cc hipMemcpy2DFromArray.cc + hipMemcpy2DFromArray_old.cc hipMemcpy2DFromArrayAsync.cc hipMemcpy2DFromArrayAsync_old.cc hipMemcpyAtoH.cc + hipMemcpyAtoH_old.cc hipMemcpyHtoA.cc + hipMemcpyHtoA_old.cc hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc @@ -112,6 +116,7 @@ set(TEST_SRC memset.cc malloc.cc hipMemcpy2DToArray.cc + hipMemcpy2DToArray_old.cc hipMemcpy2DToArrayAsync.cc hipMemcpy2DToArrayAsync_old.cc hipMemcpy3D.cc @@ -121,10 +126,13 @@ set(TEST_SRC hipMemcpy2D.cc hipMemcpy2DAsync.cc hipMemcpy2DFromArray.cc + hipMemcpy2DFromArray_old.cc hipMemcpy2DFromArrayAsync.cc hipMemcpy2DFromArrayAsync_old.cc hipMemcpyAtoH.cc + hipMemcpyAtoH_old.cc hipMemcpyHtoA.cc + hipMemcpyHtoA_old.cc hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc diff --git a/catch/unit/memory/hipMemcpy2DFromArray.cc b/catch/unit/memory/hipMemcpy2DFromArray.cc index 1972a5f10a..809e18fb6e 100644 --- a/catch/unit/memory/hipMemcpy2DFromArray.cc +++ b/catch/unit/memory/hipMemcpy2DFromArray.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -16,316 +16,239 @@ 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. */ - /* -This file verifies the following scenarios of hipMemcpy2DFromArray API -1. Negative Scenarios -2. Extent Validation Scenarios -3. hipMemcpy2DFromArray Basic Scenario -4. Pinned Memory scenarios on same and peer GPU -5. Device Context change scenario where memory is allocated in - one GPU and API is triggered from peer GPU. +Testcase Scenarios : +Unit_hipMemcpy2DFromArray_Positive_Default - Test basic memcpy between 2D array +and host/device with hipMemcpy2DFromArray api +Unit_hipMemcpy2DFromArray_Positive_Synchronization_Behavior - Test +synchronization behavior for hipMemcpy2DFromArray api +Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight - Test that no data is copied +when width/height is set to 0 Unit_hipMemcpy2DFromArray_Negative_Parameters - +Test unsuccessful execution of hipMemcpy2DFromArray api when parameters are +invalid */ +#include "array_memcpy_tests_common.hh" +#include #include -#include +#include +#include -static constexpr auto NUM_W{10}; -static constexpr auto NUM_H{10}; -/* - * This testcase verifies device to host copy for hipMemcpy2DFromArray API - * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - */ -TEST_CASE("Unit_hipMemcpy2DFromArray_Basic") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice)); +TEST_CASE("Unit_hipMemcpy2DFromArray_Positive_Default") { + using namespace std::placeholders; - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); -} - -/* - * This testcase verifies the extent validation scenarios - * of hipMemcpy2DFromArray API - */ -TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr}; - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - HipTest::initArrays(nullptr, nullptr, nullptr, - nullptr, &valData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - - SECTION("Destination width is 0") { - REQUIRE(hipMemcpy2DFromArray(A_h, 0, A_d, - 0, 0, NUM_W*sizeof(float), - NUM_H, hipMemcpyDeviceToHost) != hipSuccess); - } - // hipMemcpy2DFromArray API would return success for width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying A_d-->hData variable - // with height 0(copy will not be performed) - // 3 validating hData<-->A_h which will not be equal as copy is not done. - SECTION("Height is 0") { - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - A_h, width, width, - NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, width, - 0, hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); - } - // hipMemcpy2DFromArray API would return success for width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying A_d-->hData variable - // with width 0(copy will not be performed) - // 3 validating hData<-->A_h which will not be equal as copy is not done. - - SECTION("Width is 0") { - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - A_h, width, width, - NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, 0, - NUM_H, hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); + SECTION("Array to host") { + Memcpy2DHostFromAShell( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToHost), + width, height); } - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); - HipTest::freeArrays(nullptr, nullptr, nullptr, - nullptr, valData, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DFromArray API by copying the - * data from pinned host memory to device on same GPU - * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with PinnedMem[0](i.e., 10) - */ -TEST_CASE("Unit_hipMemcpy2DFromArray_PinnedMemSameGPU") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - constexpr auto def_val{10}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *PinnMem{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - for (int i = 0; i < NUM_W*NUM_H; i++) { - PinnMem[i] = def_val + i; + SECTION("Array to host with default kind") { + Memcpy2DHostFromAShell(std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, + width * sizeof(int), height, hipMemcpyDefault), + width, height); } - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, - width, width, NUM_H, - hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(PinnMem)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DFromArray API by copying the - * data from pinned host memory to device from Peer GPU. - * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 - * INPUT: Intializa A_d with A_h - * Copy A_d->E_h which is a pinned host memory - * OUTPUT: For validating the result,Copying A_d device variable - * --> E_h host variable - * and verifying A_h with E_h - */ -TEST_CASE("Unit_hipMemcpy2DFromArray_multiDevicePinnedMemPeerGpu") { - int numDevices = 0; - constexpr auto def_val{10}; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *E_h{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, A_h, - width, width, NUM_H, - hipMemcpyHostToDevice)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); - for (int i = 0; i < NUM_W*NUM_H; i++) { - E_h[i] = def_val + i; - } - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpy2DFromArray(E_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(E_h)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); - } else { - SUCCEED("Device Does not have P2P capability"); +#if HT_NVIDIA // EXSWHTEC-120 + SECTION("Array to device") { + SECTION("Peer access disabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToDevice), + width, height); } - } else { - SUCCEED("Number of devices are < 2"); - } -} - -/* - * This scenario verifies the hipMemcpy2DFromArray API in case of device - * context change. - * Memory is allocated in GPU-0 and the API is triggered from GPU-1 - * INPUT: Copying Host variable hData(Initial value Phi) - * --> A_d device variable - * whose memory is allocated in GPU 0 - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - * */ -TEST_CASE("Unit_hipMemcpy2DFromArray_multiDeviceContextChange") { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice)); - - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); - } else { - SUCCEED("Device Does not have P2P capability"); + SECTION("Peer access enabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToDevice), + width, height); } - } else { - SUCCEED("Number of devices are < 2"); - } -} -/* This testcase verifies the negative scenarios of - * hipMemcpy2DFromArray API - */ -TEST_CASE("Unit_hipMemcpy2DFromArray_Negative") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - - SECTION("Nullptr to destination") { - REQUIRE(hipMemcpy2DFromArray(nullptr, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost) != hipSuccess); } - SECTION("Nullptr to source") { - REQUIRE(hipMemcpy2DFromArray(A_h, width, nullptr, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost) != hipSuccess); + SECTION("Array to device with default kind") { + SECTION("Peer access disabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDefault), + width, height); + } + SECTION("Peer access enabled") { + Memcpy2DDeviceFromAShell( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDefault), + width, height); + } } - - SECTION("Passing offset more than 0") { - REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 1, - 1, width, NUM_H, - hipMemcpyDeviceToHost) != hipSuccess); - } - - SECTION("Passing array more than allocated") { - REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 0, - 0, width+2, NUM_H+2, - hipMemcpyDeviceToHost) != hipSuccess); - } - - // Cleaning of memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); +#endif } +TEST_CASE("Unit_hipMemcpy2DFromArray_Positive_Synchronization_Behavior") { + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Array to host") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyAtoHPageableSyncBehavior(std::bind(hipMemcpy2DFromArray, _1, width * sizeof(int), _2, 0, + 0, width * sizeof(int), height, hipMemcpyDeviceToHost), + width, height, true); + MemcpyAtoHPinnedSyncBehavior(std::bind(hipMemcpy2DFromArray, _1, width * sizeof(int), _2, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost), + width, height, true); + } +#if HT_NVIDIA // EXSWHTEC-214 + SECTION("Array to device") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyAtoDSyncBehavior(std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice), + width, height, false); + } +#endif +} + +TEST_CASE("Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight") { + using namespace std::placeholders; + const auto width = 16; + const auto height = 16; + + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToHost), + width, height); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToHost), + width, height); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToDevice), + width, height); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToDevice), + width, height); + } + } +} + +TEST_CASE("Unit_hipMemcpy2DFromArray_Negative_Parameters") { + using namespace std::placeholders; + + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + + const unsigned int flag = hipArrayDefault; + + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + SECTION("Array to host") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArray(nullptr, 2 * width * sizeof(int), array_alloc.ptr(), 0, + 0, width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), nullptr, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidHandle); + } +#if HT_NVIDIA // EXSWHTEC-119 + SECTION("dpitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(host_alloc.ptr(), width * sizeof(int) - 10, array_alloc.ptr(), 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 1, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 1, + width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int) + 1, height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int), height + 1, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } +#endif + } + SECTION("Array to device") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArray(nullptr, device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), nullptr, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidHandle); + } +#if HT_NVIDIA // EXSWHTEC-119 + SECTION("dpitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(device_alloc.ptr(), width * sizeof(int) - 10, array_alloc.ptr(), 0, + 0, width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 1, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 1, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int), height + 1, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } +#endif + } +} diff --git a/catch/unit/memory/hipMemcpy2DFromArray_old.cc b/catch/unit/memory/hipMemcpy2DFromArray_old.cc new file mode 100644 index 0000000000..1972a5f10a --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DFromArray_old.cc @@ -0,0 +1,331 @@ +/* +Copyright (c) 2021 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. +*/ + +/* +This file verifies the following scenarios of hipMemcpy2DFromArray API +1. Negative Scenarios +2. Extent Validation Scenarios +3. hipMemcpy2DFromArray Basic Scenario +4. Pinned Memory scenarios on same and peer GPU +5. Device Context change scenario where memory is allocated in + one GPU and API is triggered from peer GPU. +*/ + +#include +#include + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{10}; +/* + * This testcase verifies device to host copy for hipMemcpy2DFromArray API + * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_Basic") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + +/* + * This testcase verifies the extent validation scenarios + * of hipMemcpy2DFromArray API + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr}; + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + HipTest::initArrays(nullptr, nullptr, nullptr, + nullptr, &valData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + + SECTION("Destination width is 0") { + REQUIRE(hipMemcpy2DFromArray(A_h, 0, A_d, + 0, 0, NUM_W*sizeof(float), + NUM_H, hipMemcpyDeviceToHost) != hipSuccess); + } + // hipMemcpy2DFromArray API would return success for width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying A_d-->hData variable + // with height 0(copy will not be performed) + // 3 validating hData<-->A_h which will not be equal as copy is not done. + SECTION("Height is 0") { + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + A_h, width, width, + NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, + 0, 0, width, + 0, hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); + } + // hipMemcpy2DFromArray API would return success for width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying A_d-->hData variable + // with width 0(copy will not be performed) + // 3 validating hData<-->A_h which will not be equal as copy is not done. + + SECTION("Width is 0") { + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + A_h, width, width, + NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, + 0, 0, 0, + NUM_H, hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true); + } + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, + nullptr, valData, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DFromArray API by copying the + * data from pinned host memory to device on same GPU + * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with PinnedMem[0](i.e., 10) + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_PinnedMemSameGPU") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + constexpr auto def_val{10}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *PinnMem{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + for (int i = 0; i < NUM_W*NUM_H; i++) { + PinnMem[i] = def_val + i; + } + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, + width, width, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(PinnMem)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DFromArray API by copying the + * data from pinned host memory to device from Peer GPU. + * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 + * INPUT: Intializa A_d with A_h + * Copy A_d->E_h which is a pinned host memory + * OUTPUT: For validating the result,Copying A_d device variable + * --> E_h host variable + * and verifying A_h with E_h + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_multiDevicePinnedMemPeerGpu") { + int numDevices = 0; + constexpr auto def_val{10}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *E_h{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, A_h, + width, width, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); + for (int i = 0; i < NUM_W*NUM_H; i++) { + E_h[i] = def_val + i; + } + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpy2DFromArray(E_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(E_h)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); + } else { + SUCCEED("Device Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + +/* + * This scenario verifies the hipMemcpy2DFromArray API in case of device + * context change. + * Memory is allocated in GPU-0 and the API is triggered from GPU-1 + * INPUT: Copying Host variable hData(Initial value Phi) + * --> A_d device variable + * whose memory is allocated in GPU 0 + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + * */ +TEST_CASE("Unit_hipMemcpy2DFromArray_multiDeviceContextChange") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); + } else { + SUCCEED("Device Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} +/* This testcase verifies the negative scenarios of + * hipMemcpy2DFromArray API + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + + SECTION("Nullptr to destination") { + REQUIRE(hipMemcpy2DFromArray(nullptr, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost) != hipSuccess); + } + + SECTION("Nullptr to source") { + REQUIRE(hipMemcpy2DFromArray(A_h, width, nullptr, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost) != hipSuccess); + } + + SECTION("Passing offset more than 0") { + REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 1, + 1, width, NUM_H, + hipMemcpyDeviceToHost) != hipSuccess); + } + + SECTION("Passing array more than allocated") { + REQUIRE(hipMemcpy2DFromArray(A_h, width, A_d, 0, + 0, width+2, NUM_H+2, + hipMemcpyDeviceToHost) != hipSuccess); + } + + // Cleaning of memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + diff --git a/catch/unit/memory/hipMemcpy2DToArray.cc b/catch/unit/memory/hipMemcpy2DToArray.cc index e75e2c6af3..15a518ce96 100644 --- a/catch/unit/memory/hipMemcpy2DToArray.cc +++ b/catch/unit/memory/hipMemcpy2DToArray.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -16,317 +16,234 @@ 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. */ - /* -This file verifies the following scenarios of hipMemcpy2DToArray API -1. Negative Scenarios -2. Extent Validation Scenarios -3. hipMemcpy2DToArray Basic Scenario -4. Pinned Memory scenarios on same and peer GPU -5. Device Context change scenario where memory is allocated in - one GPU and API is triggered from peer GPU. +Testcase Scenarios : +Unit_hipMemcpy2DToArray_Positive_Default - Test basic memcpy between host/device +and 2D array with hipMemcpy2DToArray api +Unit_hipMemcpy2DToArray_Positive_Synchronization_Behavior - Test synchronization +behavior for hipMemcpy2DToArray api +Unit_hipMemcpy2DToArray_Positive_ZeroWidthHeight - Test that no data is copied +when width/height is set to 0 Unit_hipMemcpy2DToArray_Negative_Parameters - Test +unsuccessful execution of hipMemcpy2DToArray api when parameters are invalid */ +#include "array_memcpy_tests_common.hh" +#include #include -#include -#include +#include +#include -static constexpr auto NUM_W{10}; -static constexpr auto NUM_H{10}; -/* - * This Scenario copies the data from host to device - * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - */ -TEST_CASE("Unit_hipMemcpy2DToArray_Basic") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice)); +TEST_CASE("Unit_hipMemcpy2DToArray_Positive_Default") { + using namespace std::placeholders; - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); -} - -/* - * This testcase verifies the extent validation scenarios - */ -TEST_CASE("Unit_hipMemcpy2DToArray_ExtentValidation") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - - SECTION("Source width is 0") { - REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, hData, 0, - width, NUM_H, - hipMemcpyHostToDevice) != hipSuccess); - } - // hipMemcpy2DToArray API would return success for width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying hData(Phi)-->A_d device variable - // with height 0(copy will not be performed) - // 3.copying A_d-->hData and validating it with A_h data - SECTION("Height is 0") { - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - A_h, width, width, - NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - hData, width, - width, 0, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); - } - // hipMemcpy2DToArray API would return success for width and height as 0 - // and does not perform any copy - // Validating the result with the initialized value - // 1.Initializing A_d with Pi value - // 2.copying hData(Phi)-->A_d device variable - // with width 0(copy will not be performed) - // 3.copying A_d-->hData and validating it with A_h data - SECTION("Width is 0") { - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - A_h, width, width, - NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, - hData, width, - 0, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); + SECTION("Host to Array") { + Memcpy2DHosttoAShell(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, + width * sizeof(int), height, hipMemcpyHostToDevice), + width, height); } - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DToArray API by copying the - * data from pinned host memory to device on same GPU - * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) - * --> A_d device variable - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with PinnedMem[0](i.e., 10) - */ -TEST_CASE("Unit_hipMemcpy2DToArray_PinnedMemSameGPU") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - constexpr auto def_val{10}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *PinnMem{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - for (int i = 0; i < NUM_W*NUM_H; i++) { - PinnMem[i] = def_val + i; + SECTION("Host to Array with default kind") { + Memcpy2DHosttoAShell(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, + width * sizeof(int), height, hipMemcpyDefault), + width, height); } - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, - width, width, NUM_H, - hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(PinnMem)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); -} -/* - * This Scenario Verifies hipMemcpy2DToArray API by copying the - * data from pinned host memory to device from Peer GPU. - * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 - * INPUT: Copying Host variable E_h(Initialized with value 10+i(numelements)) - * --> A_d device variable - * whose memory is allocated in GPU 0 - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with E_h[0]+i(i.e., 10+i) - */ -TEST_CASE("Unit_hipMemcpy2DToArray_multiDevicePinnedMemPeerGpu") { - int numDevices = 0; - constexpr auto def_val{10}; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *E_h{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, nullptr, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); - for (int i = 0; i < NUM_W*NUM_H; i++) { - E_h[i] = def_val + i; - } - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, E_h, - width, width, NUM_H, - hipMemcpyHostToDevice)); - HIP_CHECK(hipSetDevice(0)); - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HIP_CHECK(hipHostFree(E_h)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, nullptr, nullptr, false); - } else { - SUCCEED("Machine Does not have P2P capability"); +#if HT_NVIDIA // EXSWHTEC-120 + SECTION("Device to Array") { + SECTION("Peer access disabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDeviceToDevice), + width, height); } - } else { - SUCCEED("Number of devices are < 2"); - } -} - -/* - * This scenario verifies the hipMemcpy2DToArray API in case of device - * context change. - * Memory is allocated in GPU-0 and the API is triggered from GPU-1 - * INPUT: Copying Host variable hData(Initial value Phi) - * --> A_d device variable - * whose memory is allocated in GPU 0 - * OUTPUT: For validating the result,Copying A_d device variable - * --> A_h host variable - * and verifying A_h with Phi - * */ -TEST_CASE("Unit_hipMemcpy2DToArray_multiDeviceDeviceContextChange") { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int canAccessPeer = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice)); - - HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost)); - REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); - - // Cleaning the memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); - } else { - SUCCEED("Machine Does not have P2P capability"); + SECTION("Peer access enabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDeviceToDevice), + width, height); } - } else { - SUCCEED("Number of devices are < 2"); - } -} -/* This testcase verifies the negative scenarios - */ -TEST_CASE("Unit_hipMemcpy2DToArray_Negative") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d{nullptr}; - size_t width{sizeof(float)*NUM_W}; - float *A_h{nullptr}, *hData{nullptr}; - - // Initialization of variables - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &hData, nullptr, - width*NUM_H, false); - HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - - SECTION("Nullptr to destination") { - REQUIRE(hipMemcpy2DToArray(nullptr, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice) != hipSuccess); } - SECTION("Nullptr to source") { - REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, - nullptr, width, width, - NUM_H, hipMemcpyHostToDevice) != hipSuccess); + SECTION("Device to Array with default kind") { + SECTION("Peer access disabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDefault), + width, height); + } + SECTION("Peer access enabled") { + Memcpy2DDevicetoAShell( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDefault), + width, height); + } } - - SECTION("Passing offset more than 0") { - REQUIRE(hipMemcpy2DToArray(A_d, 1, 1, - hData, width, width, - NUM_H, hipMemcpyHostToDevice) != hipSuccess); - } - - SECTION("Passing array more than allocated") { - REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, - hData, width, width+2, - NUM_H+2, hipMemcpyHostToDevice) != hipSuccess); - } - - // Cleaning of memory - HIP_CHECK(hipFreeArray(A_d)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, hData, nullptr, false); +#endif } +TEST_CASE("Unit_hipMemcpy2DToArray_Positive_Synchronization_Behavior") { + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Host to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyHtoASyncBehavior(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + width, height, true); + } +#if HT_NVIDIA // EXSWHTEC-214 + SECTION("Device to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + + MemcpyDtoASyncBehavior(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), + height, hipMemcpyDeviceToDevice), + width, height, false); + } +#endif +} + +TEST_CASE("Unit_hipMemcpy2DToArray_Positive_ZeroWidthHeight") { + using namespace std::placeholders; + const auto width = 16; + const auto height = 16; + + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyHostToDevice), + width, height); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyHostToDevice), width, + height); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyDeviceToDevice), + width, height); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyDeviceToDevice), + width, height); + } + } +} + +TEST_CASE("Unit_hipMemcpy2DToArray_Negative_Parameters") { + using namespace std::placeholders; + + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + + const unsigned int flag = hipArrayDefault; + + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + SECTION("Host to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArray(nullptr, 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, nullptr, 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + } +#if HT_NVIDIA // EXSWHTEC-119 + SECTION("spitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), width * sizeof(int) - 10, + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 1, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 1, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int) + 1, height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height + 1, hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } +#endif + } + SECTION("Device to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArray(nullptr, 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, nullptr, device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } +#if HT_NVIDIA // EXSWHTEC-119 + SECTION("spitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), width * sizeof(int) - 10, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 1, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 1, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height + 1, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } +#endif + } +} diff --git a/catch/unit/memory/hipMemcpy2DToArray_old.cc b/catch/unit/memory/hipMemcpy2DToArray_old.cc new file mode 100644 index 0000000000..e75e2c6af3 --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DToArray_old.cc @@ -0,0 +1,332 @@ +/* +Copyright (c) 2021 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. +*/ + +/* +This file verifies the following scenarios of hipMemcpy2DToArray API +1. Negative Scenarios +2. Extent Validation Scenarios +3. hipMemcpy2DToArray Basic Scenario +4. Pinned Memory scenarios on same and peer GPU +5. Device Context change scenario where memory is allocated in + one GPU and API is triggered from peer GPU. +*/ + +#include +#include +#include + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{10}; +/* + * This Scenario copies the data from host to device + * INPUT: Copying Host variable hData(Initialized with value Phi(1.618)) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + */ +TEST_CASE("Unit_hipMemcpy2DToArray_Basic") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + +/* + * This testcase verifies the extent validation scenarios + */ +TEST_CASE("Unit_hipMemcpy2DToArray_ExtentValidation") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + + SECTION("Source width is 0") { + REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, hData, 0, + width, NUM_H, + hipMemcpyHostToDevice) != hipSuccess); + } + // hipMemcpy2DToArray API would return success for width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying hData(Phi)-->A_d device variable + // with height 0(copy will not be performed) + // 3.copying A_d-->hData and validating it with A_h data + SECTION("Height is 0") { + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + A_h, width, width, + NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + hData, width, + width, 0, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); + } + // hipMemcpy2DToArray API would return success for width and height as 0 + // and does not perform any copy + // Validating the result with the initialized value + // 1.Initializing A_d with Pi value + // 2.copying hData(Phi)-->A_d device variable + // with width 0(copy will not be performed) + // 3.copying A_d-->hData and validating it with A_h data + SECTION("Width is 0") { + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + A_h, width, width, + NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, + hData, width, + 0, NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) == true); + } + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DToArray API by copying the + * data from pinned host memory to device on same GPU + * INPUT: Copying Host variable PinnMem(Initialized with value "10" ) + * --> A_d device variable + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with PinnedMem[0](i.e., 10) + */ +TEST_CASE("Unit_hipMemcpy2DToArray_PinnedMemSameGPU") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + constexpr auto def_val{10}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *PinnMem{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&PinnMem), width * NUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + for (int i = 0; i < NUM_W*NUM_H; i++) { + PinnMem[i] = def_val + i; + } + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, + width, width, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(PinnMem)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DToArray API by copying the + * data from pinned host memory to device from Peer GPU. + * Device Memory is allocated in GPU 0 and the API is trigerred from GPU1 + * INPUT: Copying Host variable E_h(Initialized with value 10+i(numelements)) + * --> A_d device variable + * whose memory is allocated in GPU 0 + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with E_h[0]+i(i.e., 10+i) + */ +TEST_CASE("Unit_hipMemcpy2DToArray_multiDevicePinnedMemPeerGpu") { + int numDevices = 0; + constexpr auto def_val{10}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *E_h{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, nullptr, nullptr); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&E_h), width * NUM_H)); + for (int i = 0; i < NUM_W*NUM_H; i++) { + E_h[i] = def_val + i; + } + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, E_h, + width, width, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, E_h, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(E_h)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} + +/* + * This scenario verifies the hipMemcpy2DToArray API in case of device + * context change. + * Memory is allocated in GPU-0 and the API is triggered from GPU-1 + * INPUT: Copying Host variable hData(Initial value Phi) + * --> A_d device variable + * whose memory is allocated in GPU 0 + * OUTPUT: For validating the result,Copying A_d device variable + * --> A_h host variable + * and verifying A_h with Phi + * */ +TEST_CASE("Unit_hipMemcpy2DToArray_multiDeviceDeviceContextChange") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); + } else { + SUCCEED("Machine Does not have P2P capability"); + } + } else { + SUCCEED("Number of devices are < 2"); + } +} +/* This testcase verifies the negative scenarios + */ +TEST_CASE("Unit_hipMemcpy2DToArray_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + + // Initialization of variables + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &hData, nullptr, + width*NUM_H, false); + HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + + SECTION("Nullptr to destination") { + REQUIRE(hipMemcpy2DToArray(nullptr, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice) != hipSuccess); + } + + SECTION("Nullptr to source") { + REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, + nullptr, width, width, + NUM_H, hipMemcpyHostToDevice) != hipSuccess); + } + + SECTION("Passing offset more than 0") { + REQUIRE(hipMemcpy2DToArray(A_d, 1, 1, + hData, width, width, + NUM_H, hipMemcpyHostToDevice) != hipSuccess); + } + + SECTION("Passing array more than allocated") { + REQUIRE(hipMemcpy2DToArray(A_d, 0, 0, + hData, width, width+2, + NUM_H+2, hipMemcpyHostToDevice) != hipSuccess); + } + + // Cleaning of memory + HIP_CHECK(hipFreeArray(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + diff --git a/catch/unit/memory/hipMemcpyAtoH.cc b/catch/unit/memory/hipMemcpyAtoH.cc index d8d28432e3..b97759a913 100644 --- a/catch/unit/memory/hipMemcpyAtoH.cc +++ b/catch/unit/memory/hipMemcpyAtoH.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -17,164 +17,43 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ /* - * Test Scenarios: - * 1. Perform host and Pinned host Memory - * 2. Perform bytecount 0 validation for hipMemcpyAtoH API - * 3. Allocate Memory from one GPU device and call hipMemcpyAtoH from Peer - * GPU device - * 4. Perform hipMemcpyAtoH Negative Scenarios - */ +Testcase Scenarios : +Unit_hipMemcpy2DFromArray_Positive_Default - Test basic memcpy between 2D array +and host/device with hipMemcpy2DFromArray api +Unit_hipMemcpy2DFromArray_Positive_Synchronization_Behavior - Test +synchronization behavior for hipMemcpy2DFromArray api +Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight - Test that no data is copied +when width/height is set to 0 Unit_hipMemcpy2DFromArray_Negative_Parameters - +Test unsuccessful execution of hipMemcpy2DFromArray api when parameters are +invalid +*/ +#include "array_memcpy_tests_common.hh" +#include #include -#include +#include +#include +TEST_CASE("Unit_hipMemcpyAtoH_Positive_Default") { + using namespace std::placeholders; -static constexpr auto NUM_W{10}; -static constexpr auto NUM_H{1}; -static constexpr auto copy_bytes{2}; + const auto width = GENERATE(512, 1024, 2048); + const auto allocation_size = width * sizeof(int); -/* -This testcase performs the basic and pinned host memory scenarios -of hipMemcpyAtoH API -Input: "A_d" initialized with "hData" Pi value -Output:"B_h" host variable output of hipMemcpyAtoH API - is then validated with "hData" - -The same scenario is then verified with pinned host memory -*/ - -TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_Basic", "[hipMemcpyAtoH]", - char, int, float) { - HIP_CHECK(hipSetDevice(0)); - // 1 refers to pinned host memory scenario - auto memtype_check = GENERATE(0, 1); - hipArray *A_d; - TestType *hData{nullptr}, *B_h{nullptr}; - size_t width{NUM_W * sizeof(TestType)}; - - // Initialization of data - if (memtype_check) { - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W, true); - } else { - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - } - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); - - // Performing API call - REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType)) - == hipSuccess); - - // Validating the result - REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); - - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - if (memtype_check) { - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - nullptr, true) == true); - } else { - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - nullptr, false) == true); - } + MemcpyAtoHShell(std::bind(hipMemcpyAtoH, _1, _2, 0, allocation_size), width); } -/* -This testcase performs the basic and pinned host memory scenarios -of hipMemcpyAtoH API -Memory is allocated in GPU-0 and the API is triggered from GPU-1 -Input: "A_d" initialized with "hData" Pi value -Output:"B_h" host variable output of hipMemcpyAtoH API - is then validated with "hData" -*/ -#if HT_AMD -TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_multiDevice-PeerDeviceContext", - "[hipMemcpyAtoH]", - char, int, float) { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int peerAccess = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); - if (!peerAccess) { - SUCCEED("Skipped the test as there is no peer access"); - } else { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d; - TestType *hData{nullptr}, *B_h{nullptr}; - size_t width{NUM_W * sizeof(TestType)}; +TEST_CASE("Unit_hipMemcpyAtoH_Positive_Synchronization_Behavior") { + using namespace std::placeholders; - // Initialization of data - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); + const auto width = GENERATE(512, 1024, 2048); + const auto height = 0; + const auto allocation_size = width * sizeof(int); - HIP_CHECK(hipDeviceSynchronize()); - // Changing the device context - HIP_CHECK(hipSetDevice(1)); - - // Performing API call - REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType)) - == hipSuccess); - // Validating the result - REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); - - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, - hData, B_h, - nullptr, false) == true); - } - } else { - SUCCEED("skipping the testcases as numDevices < 2"); - } -} -#endif -/* -This testcase verifies the negative scenarios of hipMemcpyAtoH API -*/ -TEST_CASE("Unit_hipMemcpyAtoH_Negative") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d; - float *hData{nullptr}, *B_h{nullptr}; - size_t width{NUM_W * sizeof(float)}; - - // Initialization of data - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); - - SECTION("Destination pointer is nullptr") { - REQUIRE(hipMemcpyAtoH(nullptr, A_d, 0, copy_bytes*sizeof(float)) - != hipSuccess); - } - - SECTION("Source offset is more than allocated size") { - REQUIRE(hipMemcpyAtoH(B_h, A_d, 100, copy_bytes*sizeof(float)) - != hipSuccess); - } - - SECTION("ByteCount is greater than allocated size") { - REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 12*sizeof(float)) != hipSuccess); - } - - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - nullptr, false) == true); + MemcpyAtoHPageableSyncBehavior(std::bind(hipMemcpyAtoH, _1, _2, 0, allocation_size), width, + height, true); + MemcpyAtoHPinnedSyncBehavior(std::bind(hipMemcpyAtoH, _1, _2, 0, allocation_size), width, height, + true); } /* @@ -183,37 +62,65 @@ Excluded the testcase for amd,as there is already a bug raised SWDEV-274683 */ #if HT_NVIDIA -TEST_CASE("Unit_hipMemcpyAtoH_SizeCheck") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d; - float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr}; - size_t width{NUM_W * sizeof(float)}; +TEST_CASE("Unit_hipMemcpyAtoH_Positive_ZeroCount") { + const auto width = 1024; + const auto height = 0; + const auto allocation_size = width * sizeof(int); - // Initialization of data - HipTest::initArrays(nullptr, nullptr, nullptr, - nullptr, &def_data, nullptr, NUM_W); - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - HipTest::setDefaultData(NUM_W, nullptr, def_data, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); + const unsigned int flag = hipArrayDefault; - SECTION("Passing 0 to copy bytes") { - REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 0) == hipSuccess); - REQUIRE(HipTest::checkArray(B_h, def_data, NUM_W, NUM_H) == true); - } + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); - SECTION(" Source Array is nullptr") { - REQUIRE(hipMemcpyAtoH(B_h, nullptr, 0, copy_bytes*sizeof(float)) - != hipSuccess); - } + int fill_value = 42; + std::fill_n(host_alloc.host_ptr(), width, fill_value); + HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), sizeof(int) * width, + sizeof(int) * width, 1, hipMemcpyHostToDevice)); + fill_value = 41; + std::fill_n(host_alloc.host_ptr(), width, fill_value); + HIP_CHECK(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, 0)); - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - def_data, false) == true); + ArrayFindIfNot(host_alloc.host_ptr(), static_cast(fill_value), width); } #endif + +TEST_CASE("Unit_hipMemcpyAtoH_Negative_Parameters") { + using namespace std::placeholders; + + const auto width = 1024; + const auto height = 0; + const auto allocation_size = width * sizeof(int); + + const unsigned int flag = hipArrayDefault; + + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpyAtoH(nullptr, array_alloc.ptr(), 0, allocation_size), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc.ptr(), nullptr, 0, allocation_size), + hipErrorInvalidValue); + } + SECTION("Offset is greater than allocated size") { + HIP_CHECK_ERROR( + hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), allocation_size + 10, allocation_size), + hipErrorInvalidValue); + } + SECTION("Count is greater than allocated size") { + HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, allocation_size + 10), + hipErrorInvalidValue); + } + SECTION("2D array is allocated") { + const auto width_2d = 32; + const auto height_2d = width_2d; + const auto allocation_size_2d = width_2d * height_2d * sizeof(int); + + ArrayAllocGuard array_alloc_2d(make_hipExtent(width_2d, height_2d, 0), flag); + LinearAllocGuard host_alloc_2d(LinearAllocs::hipHostMalloc, allocation_size_2d); + HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc_2d.ptr(), array_alloc_2d.ptr(), 0, allocation_size_2d), + hipErrorInvalidValue); + } +} diff --git a/catch/unit/memory/hipMemcpyAtoH_old.cc b/catch/unit/memory/hipMemcpyAtoH_old.cc new file mode 100644 index 0000000000..d8d28432e3 --- /dev/null +++ b/catch/unit/memory/hipMemcpyAtoH_old.cc @@ -0,0 +1,219 @@ +/* +Copyright (c) 2021 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. +*/ +/* + * Test Scenarios: + * 1. Perform host and Pinned host Memory + * 2. Perform bytecount 0 validation for hipMemcpyAtoH API + * 3. Allocate Memory from one GPU device and call hipMemcpyAtoH from Peer + * GPU device + * 4. Perform hipMemcpyAtoH Negative Scenarios + */ + +#include +#include + + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{1}; +static constexpr auto copy_bytes{2}; + +/* +This testcase performs the basic and pinned host memory scenarios +of hipMemcpyAtoH API +Input: "A_d" initialized with "hData" Pi value +Output:"B_h" host variable output of hipMemcpyAtoH API + is then validated with "hData" + +The same scenario is then verified with pinned host memory +*/ + +TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_Basic", "[hipMemcpyAtoH]", + char, int, float) { + HIP_CHECK(hipSetDevice(0)); + // 1 refers to pinned host memory scenario + auto memtype_check = GENERATE(0, 1); + hipArray *A_d; + TestType *hData{nullptr}, *B_h{nullptr}; + size_t width{NUM_W * sizeof(TestType)}; + + // Initialization of data + if (memtype_check) { + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W, true); + } else { + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + } + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + // Performing API call + REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType)) + == hipSuccess); + + // Validating the result + REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + if (memtype_check) { + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + nullptr, true) == true); + } else { + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + nullptr, false) == true); + } +} + +/* +This testcase performs the basic and pinned host memory scenarios +of hipMemcpyAtoH API +Memory is allocated in GPU-0 and the API is triggered from GPU-1 +Input: "A_d" initialized with "hData" Pi value +Output:"B_h" host variable output of hipMemcpyAtoH API + is then validated with "hData" +*/ +#if HT_AMD +TEMPLATE_TEST_CASE("Unit_hipMemcpyAtoH_multiDevice-PeerDeviceContext", + "[hipMemcpyAtoH]", + char, int, float) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d; + TestType *hData{nullptr}, *B_h{nullptr}; + size_t width{NUM_W * sizeof(TestType)}; + + // Initialization of data + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + HIP_CHECK(hipDeviceSynchronize()); + // Changing the device context + HIP_CHECK(hipSetDevice(1)); + + // Performing API call + REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, copy_bytes*sizeof(TestType)) + == hipSuccess); + // Validating the result + REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, + hData, B_h, + nullptr, false) == true); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} +#endif +/* +This testcase verifies the negative scenarios of hipMemcpyAtoH API +*/ +TEST_CASE("Unit_hipMemcpyAtoH_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d; + float *hData{nullptr}, *B_h{nullptr}; + size_t width{NUM_W * sizeof(float)}; + + // Initialization of data + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + SECTION("Destination pointer is nullptr") { + REQUIRE(hipMemcpyAtoH(nullptr, A_d, 0, copy_bytes*sizeof(float)) + != hipSuccess); + } + + SECTION("Source offset is more than allocated size") { + REQUIRE(hipMemcpyAtoH(B_h, A_d, 100, copy_bytes*sizeof(float)) + != hipSuccess); + } + + SECTION("ByteCount is greater than allocated size") { + REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 12*sizeof(float)) != hipSuccess); + } + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + nullptr, false) == true); +} + +/* +This testcase verifies size 0 check of hipMemcpyAtoH API +Excluded the testcase for amd,as there is already a bug raised +SWDEV-274683 +*/ +#if HT_NVIDIA +TEST_CASE("Unit_hipMemcpyAtoH_SizeCheck") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d; + float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr}; + size_t width{NUM_W * sizeof(float)}; + + // Initialization of data + HipTest::initArrays(nullptr, nullptr, nullptr, + nullptr, &def_data, nullptr, NUM_W); + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + HipTest::setDefaultData(NUM_W, nullptr, def_data, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + SECTION("Passing 0 to copy bytes") { + REQUIRE(hipMemcpyAtoH(B_h, A_d, 0, 0) == hipSuccess); + REQUIRE(HipTest::checkArray(B_h, def_data, NUM_W, NUM_H) == true); + } + + SECTION(" Source Array is nullptr") { + REQUIRE(hipMemcpyAtoH(B_h, nullptr, 0, copy_bytes*sizeof(float)) + != hipSuccess); + } + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + def_data, false) == true); +} +#endif diff --git a/catch/unit/memory/hipMemcpyHtoA.cc b/catch/unit/memory/hipMemcpyHtoA.cc index 8712704400..e144e261f6 100644 --- a/catch/unit/memory/hipMemcpyHtoA.cc +++ b/catch/unit/memory/hipMemcpyHtoA.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -16,171 +16,41 @@ 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. */ - /* - * Test Scenarios: - * 1. Perform simple and pinned host memory of hipMemcpyHtoA API - * 2. Allocate Memory from one GPU device and call hipMemcpyHtoA from Peer - * GPU device - * 3. Perform hipMemcpyHtoA Negative Scenarios - * 4. Perform bytecount 0 validation for hipMemcpyHtoA API +Testcase Scenarios : +Unit_hipMemcpyHtoA_Positive_Default - Test basic memcpy between host and 1D +array with hipMemcpyHtoA api +Unit_hipMemcpyHtoA_Positive_Synchronization_Behavior - Test synchronization +behavior for hipMemcpyHtoA api Unit_hipMemcpyHtoA_Positive_ZeroCount - Test that +no data is copied when allocation_size is set to 0 +Unit_hipMemcpyHtoA_Negative_Parameters - Test unsuccessful execution of +hipMemcpyHtoA api when parameters are invalid */ +#include "array_memcpy_tests_common.hh" +#include #include -#include +#include +#include -static constexpr auto NUM_W{10}; -static constexpr auto NUM_H{1}; -static constexpr auto copy_bytes{2}; +TEST_CASE("Unit_hipMemcpyHtoA_Positive_Default") { + using namespace std::placeholders; -/* -This testcase performs the basic and pinned host memory scenarios -of hipMemcpyHtoA API -Input: "B_h" which is initialized with 1.6 -Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable - validated the result with "B_h" + const auto width = GENERATE(512, 1024, 2048); + const auto allocation_size = width * sizeof(int); -The same scenario is then verified with pinned host memory -*/ -TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_Basic", "[hipMemcpyHtoA]", - char, int, float) { - HIP_CHECK(hipSetDevice(0)); - auto memtype_check = GENERATE(0, 1); - hipArray *A_d; - TestType *hData{nullptr}, *B_h{nullptr}; - size_t width{NUM_W * sizeof(TestType)}; - - // Initialization of data - if (memtype_check) { - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W, true); - } else { - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - } - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); - - // Performing API call - HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType))); - HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d, - 0, 0, sizeof(TestType)*NUM_W, 1, hipMemcpyDeviceToHost)); - - - // Validating the result - REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); - - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - if (memtype_check) { - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - nullptr, true) == true); - } else { - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - nullptr, false) == true); - } + MemcpyHtoAShell(std::bind(hipMemcpyHtoA, _1, 0, _2, allocation_size), width); } +TEST_CASE("Unit_hipMemcpyHtoA_Positive_Synchronization_Behavior") { + using namespace std::placeholders; -/* -This testcase performs the peer device context scenario -of hipMemcpyHtoA API -Memory is allocated in GPU-0 and the API is triggered from GPU-1 -Input: "B_h" which is initialized with 1.6 -Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable - validated the result with "B_h" -*/ -#if HT_AMD -TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_multiDevice-PeerDeviceContext", - "[hipMemcpyHtoA]", - char, int, float) { - int numDevices = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - int peerAccess = 0; - HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); - if (!peerAccess) { - SUCCEED("Skipped the test as there is no peer access"); - } else { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d; - TestType *hData{nullptr}, *B_h{nullptr}; - size_t width{NUM_W * sizeof(TestType)}; + const auto width = GENERATE(512, 1024, 2048); + const auto height = 0; + const auto allocation_size = width * sizeof(int); - // Initialization of data - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); - - // Changing the device context - HIP_CHECK(hipSetDevice(1)); - - // Performing API call - HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType))); - HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d, - 0, 0, sizeof(TestType)*NUM_W, 1, - hipMemcpyDeviceToHost)); - - // Validating the result - REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); - - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, - hData, B_h, - nullptr, false) == true); - } - } else { - SUCCEED("skipping the testcases as numDevices < 2"); - } -} -#endif - - -/* -This testcase verifies the negative scenarios of hipMemcpyHtoA API -*/ -TEST_CASE("Unit_hipMemcpyHtoA_Negative") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d; - float *hData{nullptr}, *B_h{nullptr}; - size_t width{NUM_W * sizeof(float)}; - - // Initialization of data - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); - - SECTION("Source pointer is nullptr") { - REQUIRE(hipMemcpyHtoA(A_d, 0, nullptr, copy_bytes*sizeof(float)) - != hipSuccess); - } - - SECTION("Source offset is more than allocated size") { - REQUIRE(hipMemcpyHtoA(A_d, 100, B_h, copy_bytes*sizeof(float)) - != hipSuccess); - } - - SECTION("ByteCount is greater than allocated size") { - REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 12*sizeof(float)) != hipSuccess); - } - - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - nullptr, false) == true); + MemcpyHtoASyncBehavior(std::bind(hipMemcpyHtoA, _1, 0, _2, allocation_size), width, height, true); } /* @@ -189,41 +59,69 @@ This is excluded for AMD as we have a bug already raised SWDEV-274683 */ #if HT_NVIDIA -TEST_CASE("Unit_hipMemcpyHtoA_SizeCheck") { - HIP_CHECK(hipSetDevice(0)); - hipArray *A_d; - float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr}; - size_t width{NUM_W * sizeof(float)}; +TEST_CASE("Unit_hipMemcpyHtoA_Positive_ZeroCount") { + const auto width = 1024; + const auto height = 0; + const auto allocation_size = width * sizeof(int); - // Initialization of data - HipTest::initArrays(nullptr, nullptr, nullptr, - nullptr, &def_data, nullptr, NUM_W); - HipTest::initArrays(nullptr, nullptr, nullptr, - &hData, &B_h, nullptr, NUM_W); - HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); - HipTest::setDefaultData(NUM_W, nullptr, def_data, nullptr); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); - HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, - width, NUM_H, hipMemcpyHostToDevice)); + const unsigned int flag = hipArrayDefault; - SECTION("Passing 0 to copy bytes") { - REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 0) == hipSuccess); - HIP_CHECK(hipMemcpy2DFromArray(def_data, sizeof(float)*NUM_W, A_d, - 0, 0, sizeof(float)*NUM_W, 1, - hipMemcpyDeviceToHost)); + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); - REQUIRE(HipTest::checkArray(hData, def_data, NUM_W, NUM_H) == true); - } + int fill_value = 42; + std::fill_n(host_alloc.host_ptr(), width, fill_value); + HIP_CHECK(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.host_ptr(), sizeof(int) * width, + sizeof(int) * width, 1, hipMemcpyHostToDevice)); + fill_value = 41; + std::fill_n(host_alloc.host_ptr(), width, fill_value); + HIP_CHECK(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), 0)); - SECTION(" Source Array is nullptr") { - REQUIRE(hipMemcpyHtoA(nullptr, 0, B_h, copy_bytes*sizeof(float)) - != hipSuccess); - } + HIP_CHECK(hipMemcpy2DFromArray(host_alloc.host_ptr(), sizeof(int) * width, array_alloc.ptr(), 0, + 0, sizeof(int) * width, 1, hipMemcpyDeviceToHost)); - // DeAllocating the memory - HIP_CHECK(hipFreeArray(A_d)); - REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, - def_data, false) == true); + ArrayFindIfNot(host_alloc.host_ptr(), static_cast(42), width); } #endif + +TEST_CASE("Unit_hipMemcpyHtoA_Negative_Parameters") { + using namespace std::placeholders; + + const auto width = 1024; + const auto height = 0; + const auto allocation_size = width * sizeof(int); + + const unsigned int flag = hipArrayDefault; + + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpyHtoA(nullptr, 0, host_alloc.ptr(), allocation_size), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc.ptr(), 0, nullptr, allocation_size), + hipErrorInvalidValue); + } + SECTION("Offset is greater than allocated size") { + HIP_CHECK_ERROR( + hipMemcpyHtoA(array_alloc.ptr(), allocation_size + 10, host_alloc.ptr(), allocation_size), + hipErrorInvalidValue); + } + SECTION("Count is greater than allocated size") { + HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), allocation_size + 10), + hipErrorInvalidValue); + } + + SECTION("2D array is allocated") { + const auto width_2d = 32; + const auto height_2d = width_2d; + const auto allocation_size_2d = width_2d * height_2d * sizeof(int); + + ArrayAllocGuard array_alloc_2d(make_hipExtent(width_2d, height_2d, 0), flag); + LinearAllocGuard host_alloc_2d(LinearAllocs::hipHostMalloc, allocation_size_2d); + HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc_2d.ptr(), 0, host_alloc_2d.ptr(), allocation_size_2d), + hipErrorInvalidValue); + } +} diff --git a/catch/unit/memory/hipMemcpyHtoA_old.cc b/catch/unit/memory/hipMemcpyHtoA_old.cc new file mode 100644 index 0000000000..8712704400 --- /dev/null +++ b/catch/unit/memory/hipMemcpyHtoA_old.cc @@ -0,0 +1,229 @@ +/* +Copyright (c) 2021 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. +*/ + +/* + * Test Scenarios: + * 1. Perform simple and pinned host memory of hipMemcpyHtoA API + * 2. Allocate Memory from one GPU device and call hipMemcpyHtoA from Peer + * GPU device + * 3. Perform hipMemcpyHtoA Negative Scenarios + * 4. Perform bytecount 0 validation for hipMemcpyHtoA API +*/ + +#include +#include + + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{1}; +static constexpr auto copy_bytes{2}; + +/* +This testcase performs the basic and pinned host memory scenarios +of hipMemcpyHtoA API +Input: "B_h" which is initialized with 1.6 +Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable + validated the result with "B_h" + +The same scenario is then verified with pinned host memory +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_Basic", "[hipMemcpyHtoA]", + char, int, float) { + HIP_CHECK(hipSetDevice(0)); + auto memtype_check = GENERATE(0, 1); + hipArray *A_d; + TestType *hData{nullptr}, *B_h{nullptr}; + size_t width{NUM_W * sizeof(TestType)}; + + // Initialization of data + if (memtype_check) { + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W, true); + } else { + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + } + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + // Performing API call + HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType))); + HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d, + 0, 0, sizeof(TestType)*NUM_W, 1, hipMemcpyDeviceToHost)); + + + // Validating the result + REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + if (memtype_check) { + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + nullptr, true) == true); + } else { + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + nullptr, false) == true); + } +} + + +/* +This testcase performs the peer device context scenario +of hipMemcpyHtoA API +Memory is allocated in GPU-0 and the API is triggered from GPU-1 +Input: "B_h" which is initialized with 1.6 +Output: "A_d" output of hipMemcpyHtoA is copied to "hData" host variable + validated the result with "B_h" +*/ +#if HT_AMD +TEMPLATE_TEST_CASE("Unit_hipMemcpyHtoA_multiDevice-PeerDeviceContext", + "[hipMemcpyHtoA]", + char, int, float) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d; + TestType *hData{nullptr}, *B_h{nullptr}; + size_t width{NUM_W * sizeof(TestType)}; + + // Initialization of data + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + // Changing the device context + HIP_CHECK(hipSetDevice(1)); + + // Performing API call + HIP_CHECK(hipMemcpyHtoA(A_d, 0, B_h, copy_bytes*sizeof(TestType))); + HIP_CHECK(hipMemcpy2DFromArray(hData, sizeof(TestType)*NUM_W, A_d, + 0, 0, sizeof(TestType)*NUM_W, 1, + hipMemcpyDeviceToHost)); + + // Validating the result + REQUIRE(HipTest::checkArray(B_h, hData, copy_bytes, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, + hData, B_h, + nullptr, false) == true); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} +#endif + + +/* +This testcase verifies the negative scenarios of hipMemcpyHtoA API +*/ +TEST_CASE("Unit_hipMemcpyHtoA_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d; + float *hData{nullptr}, *B_h{nullptr}; + size_t width{NUM_W * sizeof(float)}; + + // Initialization of data + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + SECTION("Source pointer is nullptr") { + REQUIRE(hipMemcpyHtoA(A_d, 0, nullptr, copy_bytes*sizeof(float)) + != hipSuccess); + } + + SECTION("Source offset is more than allocated size") { + REQUIRE(hipMemcpyHtoA(A_d, 100, B_h, copy_bytes*sizeof(float)) + != hipSuccess); + } + + SECTION("ByteCount is greater than allocated size") { + REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 12*sizeof(float)) != hipSuccess); + } + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + nullptr, false) == true); +} + +/* +This testcase verifies the size 0 check of hipMemcpyHtoA API +This is excluded for AMD as we have a bug already raised +SWDEV-274683 +*/ +#if HT_NVIDIA +TEST_CASE("Unit_hipMemcpyHtoA_SizeCheck") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d; + float *hData{nullptr}, *B_h{nullptr}, *def_data{nullptr}; + size_t width{NUM_W * sizeof(float)}; + + // Initialization of data + HipTest::initArrays(nullptr, nullptr, nullptr, + nullptr, &def_data, nullptr, NUM_W); + HipTest::initArrays(nullptr, nullptr, nullptr, + &hData, &B_h, nullptr, NUM_W); + HipTest::setDefaultData(NUM_W, hData, B_h, nullptr); + HipTest::setDefaultData(NUM_W, nullptr, def_data, nullptr); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice)); + + SECTION("Passing 0 to copy bytes") { + REQUIRE(hipMemcpyHtoA(A_d, 0, B_h, 0) == hipSuccess); + HIP_CHECK(hipMemcpy2DFromArray(def_data, sizeof(float)*NUM_W, A_d, + 0, 0, sizeof(float)*NUM_W, 1, + hipMemcpyDeviceToHost)); + + REQUIRE(HipTest::checkArray(hData, def_data, NUM_W, NUM_H) == true); + } + + SECTION(" Source Array is nullptr") { + REQUIRE(hipMemcpyHtoA(nullptr, 0, B_h, copy_bytes*sizeof(float)) + != hipSuccess); + } + + // DeAllocating the memory + HIP_CHECK(hipFreeArray(A_d)); + REQUIRE(HipTest::freeArrays(nullptr, nullptr, nullptr, hData, B_h, + def_data, false) == true); +} +#endif