diff --git a/projects/hip/tests/catch/include/hip_test_checkers.hh b/projects/hip/tests/catch/include/hip_test_checkers.hh index 4b2e2f49ba..a96e597b37 100644 --- a/projects/hip/tests/catch/include/hip_test_checkers.hh +++ b/projects/hip/tests/catch/include/hip_test_checkers.hh @@ -1,5 +1,15 @@ #pragma once #include "hip_test_common.hh" +#include +using namespace std; +#define guarantee(cond, str) \ + { \ + if (!(cond)) { \ + std::cout << str << std::endl; \ + abort(); \ + } \ + } + namespace HipTest { template @@ -38,6 +48,34 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM return mismatchCount; } +template // pointer type +void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) { + for (size_t i = 0; i < depth; i++) { + for (size_t j = 0; j < height; j++) { + for (size_t k = 0; k < width; k++) { + int offset = i*width*height + j*width + k; + if (hData[offset] != hOutputData[offset]) { + cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" + << hOutputData[offset]<<" "; + cout << "mistmatch at: " << i<< j< // pointer type +bool checkArray(T *result, T *compare, size_t width, size_t height) { + for (size_t i = 0; i < height; i++) { + for (size_t j = 0; j < width; j++) { + if (result[(i*width) + j] != compare[(i*width) + j]) { + std::cout << result[(i*width) + j] << "\t" << compare[(i*width) + j] << std::endl; + return false; + } + } + } + return true; +} template size_t checkVectorADD(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch = true, @@ -62,10 +100,21 @@ void checkTest(T* expected_H, T* result_H, size_t N, bool expectMatch = true) { template void setDefaultData(size_t numElements, T* A_h, T* B_h, T* C_h) { // Initialize the host data: + for (size_t i = 0; i < numElements; i++) { - if (A_h) (A_h)[i] = 3.146f + i; // Pi - if (B_h) (B_h)[i] = 1.618f + i; // Phi - if (C_h) (C_h)[i] = 0.0f + i; + if (std::is_same::value || std::is_same::value) { + if (A_h) (A_h)[i] = 3; + if (B_h) (B_h)[i] = 4; + if (C_h) (C_h)[i] = 5; + } else if(std::is_same::value || std::is_same::value) { + if (A_h) (A_h)[i] = 'a'; + if (B_h) (B_h)[i] = 'b'; + if (C_h) (C_h)[i] = 'c'; + } else { + if (A_h) (A_h)[i] = 3.146f + i; + if (B_h) (B_h)[i] = 1.618f + i; + if (C_h) (C_h)[i] = 1.4f + i; + } } } diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index ac75271618..e2593351a1 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -2,6 +2,8 @@ set(TEST_SRC memset.cc malloc.cc + hipMemcpy2DToArray.cc + hipMemcpy2DToArrayAsync.cc ) # Create shared lib of all tests diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DToArray.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DToArray.cc new file mode 100644 index 0000000000..a58f46f417 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DToArray.cc @@ -0,0 +1,332 @@ +/* +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 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; + 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; + 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/projects/hip/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc new file mode 100644 index 0000000000..1439f9a003 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc @@ -0,0 +1,362 @@ +/* +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 hipMemcpy2DToArrayAsync API +1. Negative Scenarios +2. Extent Validation Scenarios +3. hipMemcpy2DToArrayAsync 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 +#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_hipMemcpy2DToArrayAsync_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(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} + +/* + * This testcase verifies the extent validation scenarios + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_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("Source width is 0") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, 0, + width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + } + // 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(hipMemcpy2DToArrayAsync(A_d, 0, 0, A_h, width, + width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, 0, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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(hipMemcpy2DToArrayAsync(A_d, 0, 0, A_h, width, + width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + 0, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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)); + HIP_CHECK(hipStreamDestroy(stream)); + 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_hipMemcpy2DToArrayAsync_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(hipMemcpy2DToArrayAsync(A_d, 0, 0, PinnMem, + width, width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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)); + HIP_CHECK(hipStreamDestroy(stream)); + 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_hipMemcpy2DToArrayAsync_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(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(hipMemcpy2DToArrayAsync(A_d, 0, 0, E_h, width, + width, NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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)); + HIP_CHECK(hipStreamDestroy(stream)); + 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_hipMemcpy2DToArrayAsync_multiDeviceDeviceContextChange") { + 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(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, width, + NUM_H, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + 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)); + HIP_CHECK(hipStreamDestroy(stream)); + 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_hipMemcpy2DToArrayAsync_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(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, width, + width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + SECTION("Nullptr to source") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, nullptr, + width, width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + SECTION("Passing offset more than 0") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 1, 1, hData, width, + width, NUM_H, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + SECTION("Passing array more than allocated") { + REQUIRE(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width+2, NUM_H+2, hipMemcpyHostToDevice, + stream) != hipSuccess); + } + + // Cleaning of Memory + HIP_CHECK(hipFreeArray(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, hData, nullptr, false); +} +