From 63e3233c20c5d8e8cd3fae51916dd919fc58c19f Mon Sep 17 00:00:00 2001 From: DURGESH KROTTAPALLI Date: Mon, 28 Jun 2021 15:55:07 +0530 Subject: [PATCH] SWDEV-292393 - [catch2][dtest] hipMemcpyAtoH and hipMemcpyHtoA APIs to catch2 Migrated functional and negative scenarios of hipMemcpyAtoH and hipMemcpyHtoA APIs to catch2 framework Change-Id: I68bb37d99fc371b8803e64ebf1533c0870b14fab [ROCm/hip commit: bdc90769fb8478165d744d209c76735ce3691814] --- .../tests/catch/unit/memory/CMakeLists.txt | 2 + .../tests/catch/unit/memory/hipMemcpyAtoH.cc | 219 +++++++++++++++++ .../tests/catch/unit/memory/hipMemcpyHtoA.cc | 229 ++++++++++++++++++ 3 files changed, 450 insertions(+) create mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyAtoH.cc create mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyHtoA.cc diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index 773f75f71f..b82efb6b32 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -14,6 +14,8 @@ set(TEST_SRC hipMemcpy2DAsync.cc hipMemcpy2DFromArray.cc hipMemcpy2DFromArrayAsync.cc + hipMemcpyAtoH.cc + hipMemcpyHtoA.cc ) # Create shared lib of all tests diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyAtoH.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyAtoH.cc new file mode 100644 index 0000000000..59d17afc6d --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyAtoH.cc @@ -0,0 +1,219 @@ +/* +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. +*/ +/* + * 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/projects/hip/tests/catch/unit/memory/hipMemcpyHtoA.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyHtoA.cc new file mode 100644 index 0000000000..0b2064d90e --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyHtoA.cc @@ -0,0 +1,229 @@ +/* +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. +*/ + +/* + * 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