From f6e9d2491bf4a9804dbaf0f7358fc5f06e01ce7a Mon Sep 17 00:00:00 2001 From: DURGESH KROTTAPALLI Date: Tue, 29 Jun 2021 00:39:18 +0530 Subject: [PATCH] SWDEV-292393 - [catch2][dtest] Tests for hipMemcpy2DFromArray and hipMemcpy2DFromArrayAsync apis. Added functional and negative scenarios for hipMemcpy2DFromArray and hipMemcpy2DFromArrayAsync APIs Change-Id: I269499b1d538affd505d9f529e7516cb3aa17863 --- tests/catch/unit/memory/CMakeLists.txt | 2 + .../catch/unit/memory/hipMemcpy2DFromArray.cc | 326 ++++++++++++++++ .../unit/memory/hipMemcpy2DFromArrayAsync.cc | 359 ++++++++++++++++++ 3 files changed, 687 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemcpy2DFromArray.cc create mode 100644 tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 2c90cbacd7..773f75f71f 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -12,6 +12,8 @@ set(TEST_SRC hipMemcpyParam2DAsync.cc hipMemcpy2D.cc hipMemcpy2DAsync.cc + hipMemcpy2DFromArray.cc + hipMemcpy2DFromArrayAsync.cc ) # Create shared lib of all tests diff --git a/tests/catch/unit/memory/hipMemcpy2DFromArray.cc b/tests/catch/unit/memory/hipMemcpy2DFromArray.cc new file mode 100644 index 0000000000..830c99e0cf --- /dev/null +++ b/tests/catch/unit/memory/hipMemcpy2DFromArray.cc @@ -0,0 +1,326 @@ +/* +Copyright (c) 2021 - present 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}; + // 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("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, NUM_W*sizeof(float), + 0, hipMemcpyDeviceToHost)); + REQUIRE(HipTest::checkArray(hData, A_h, 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, 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 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; + 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; + 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/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc new file mode 100644 index 0000000000..7057466937 --- /dev/null +++ b/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc @@ -0,0 +1,359 @@ +/* +Copyright (c) 2021 - present 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 hipMemcpy2DFromArrayAsync API +1. Negative Scenarios +2. Extent Validation Scenarios +3. hipMemcpy2DFromArrayAsync Basic Scenario +4. Pinned Memory scenarios on same and peer GPU +5. Device Context change scenario where memory is allocated in + one GPU and stream is created in peer GPU. +*/ + +#include +#include + + +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{10}; + +/* + * This testcase copies the data from host to device of + hipMemcpy2DFromArrayAsync 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_hipMemcpy2DFromArrayAsync_Basic") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // 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(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + +/* + * This testcase verifies the extent validation scenarios + * of hipMemcpy2DFromArrayAsync API + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + + // 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)); + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Destination width is 0") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, 0, A_d, + 0, 0, NUM_W*sizeof(float), + NUM_H, hipMemcpyDeviceToHost, stream) + != hipSuccess); + } + // hipMemcpy2DFromArrayAsync 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(hipMemcpy2DFromArrayAsync(hData, width, A_d, + 0, 0, NUM_W*sizeof(float), + 0, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true); + } + // hipMemcpy2DFromArrayAsync 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(hipMemcpy2DFromArrayAsync(hData, width, A_d, + 0, 0, 0, + NUM_H, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true); + } + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DFromArrayAsync 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_hipMemcpy2DFromArrayAsync_PinnedHostMemSameGpu") { + 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}; + hipStream_t stream; + + // 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(hipStreamCreate(&stream)); + + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, PinnMem, + width, width, NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, PinnMem, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipHostFree(PinnMem)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, nullptr, false); +} +/* + * This Scenario Verifies hipMemcpy2DFromArrayAsync 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: Initialize data, A_h --> A_d device variable + * whose memory is allocated in GPU 0 + then A_d-->E_h in GPU1 + * OUTPUT: validating the result by comparing A_h and E_h + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDevicePinnedHostMem") { + int numDevices = 0; + constexpr auto def_val{10}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + 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}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // 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(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, A_h, width, + width, NUM_H, hipMemcpyHostToDevice)); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipMemcpy2DFromArrayAsync(E_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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)); + HIP_CHECK(hipStreamDestroy(stream)); + 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 hipMemcpy2DFromArrayAsync 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_hipMemcpy2DFromArrayAsync_multiDeviceContextChange") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + int canAccessPeer = 0; + 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}; + hipStream_t stream; + + // 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(hipStreamCreate(&stream)); + HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, width, + NUM_H, hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); + + // Cleaning the memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + 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 hipMemcpy2DFromArrayAsync API + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Negative") { + HIP_CHECK(hipSetDevice(0)); + hipArray *A_d{nullptr}; + size_t width{sizeof(float)*NUM_W}; + float *A_h{nullptr}, *hData{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // 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(hipMemcpy2DFromArrayAsync(nullptr, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + SECTION("Nullptr to source") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, nullptr, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + SECTION("Passing offset more than 0") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, A_d, 1, + 1, width, NUM_H, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + SECTION("Passing array more than allocated") { + REQUIRE(hipMemcpy2DFromArrayAsync(A_h, width, A_d, 0, + 0, width+2, NUM_H+2, + hipMemcpyDeviceToHost, + stream) != hipSuccess); + } + + // Cleaning of Memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} +