From b49e8e9fdfecfaaa292bd0f188b8adc167d6bfad Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 17 May 2022 15:01:15 +0530 Subject: [PATCH] SWDEV-228443 - Enhancing hip unit tests for Memory Allocation APIs (#2616) Adding new testcases for hipMalloc3D,hipMalloc3DArray, hipArrayCreate,hipMallocPitch and hipMallocArray APIs Change-Id: Ia2cc8865d605272995aaf703dd26954d11ded2ea --- catch/include/hip_test_common.hh | 22 +- catch/unit/memory/CMakeLists.txt | 10 + catch/unit/memory/hipArrayCreate.cc | 148 +++++++++++++ catch/unit/memory/hipMalloc3D.cc | 133 ++++++++++++ catch/unit/memory/hipMalloc3DArray.cc | 195 +++++++++++++++++ catch/unit/memory/hipMallocArray.cc | 172 +++++++++++++++ catch/unit/memory/hipMallocPitch.cc | 296 ++++++++++++++++++++++++++ 7 files changed, 975 insertions(+), 1 deletion(-) create mode 100644 catch/unit/memory/hipArrayCreate.cc create mode 100644 catch/unit/memory/hipMalloc3D.cc create mode 100644 catch/unit/memory/hipMalloc3DArray.cc create mode 100644 catch/unit/memory/hipMallocArray.cc create mode 100644 catch/unit/memory/hipMallocPitch.cc diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index e767f07387..f7dc215478 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -71,7 +71,27 @@ THE SOFTWARE. printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ abort(); \ } - +#if HT_NVIDIA +#define CTX_CREATE() \ + hipCtx_t context;\ + initHipCtx(&context); +#define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context)); +#define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array)); +#define HIP_TEX_REFERENCE hipTexRef +#define HIP_ARRAY hiparray +static void initHipCtx(hipCtx_t *pcontext) { + HIPCHECK(hipInit(0)); + hipDevice_t device; + HIPCHECK(hipDeviceGet(&device, 0)); + HIPCHECK(hipCtxCreate(pcontext, 0, device)); +} +#else +#define CTX_CREATE() +#define CTX_DESTROY() +#define ARRAY_DESTROY(array) HIPCHECK(hipFreeArray(array)); +#define HIP_TEX_REFERENCE textureReference* +#define HIP_ARRAY hipArray* +#endif // Utility Functions diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 64fa0bcf62..c2355661aa 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -75,6 +75,11 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMallocPitch.cc + hipMallocArray.cc + hipMalloc3D.cc + hipMalloc3DArray.cc + hipArrayCreate.cc ) else() set(TEST_SRC @@ -129,6 +134,11 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMallocPitch.cc + hipMallocArray.cc + hipMalloc3D.cc + hipMalloc3DArray.cc + hipArrayCreate.cc ) endif() diff --git a/catch/unit/memory/hipArrayCreate.cc b/catch/unit/memory/hipArrayCreate.cc new file mode 100644 index 0000000000..3bbf6cbd54 --- /dev/null +++ b/catch/unit/memory/hipArrayCreate.cc @@ -0,0 +1,148 @@ +/* +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 +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. +*/ + +/* +hipArrayCreate API test scenarios +1. Negative Scenarios +2. Allocating Small and big chunk data +3. Multithreaded scenario +*/ + +#include + +static constexpr auto NUM_W{4}; +static constexpr auto BIGNUM_W{100}; +static constexpr auto NUM_H{4}; +static constexpr auto BIGNUM_H{100}; +static constexpr auto ARRAY_LOOP{100}; + +/* + * This API verifies memory allocations for small and + * bigger chunks of data. + * Two scenarios are verified in this API + * 1. SmallArray: Allocates NUM_W*NUM_H in a loop and + * releases the memory and verifies the meminfo. + * 2. BigArray: Allocates BIGNUM_W*BIGNUM_H in a loop and + * releases the memory and verifies the meminfo + * + * In both cases, the memory info before allocation and + * after releasing the memory should be the same. + * + */ + +static void ArrayCreate_DiffSizes(int gpu) { + HIP_CHECK(hipSetDevice(gpu)); + std::vector array_size; + array_size.push_back(NUM_W); + array_size.push_back(BIGNUM_W); + for (auto &size : array_size) { + HIP_ARRAY array[ARRAY_LOOP]; + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_ARRAY_DESCRIPTOR desc; + desc.NumChannels = 1; + if (size == NUM_W) { + desc.Width = NUM_W; + desc.Height = NUM_H; + } else { + desc.Width = BIGNUM_W; + desc.Height = BIGNUM_H; + } + desc.Format = HIP_AD_FORMAT_FLOAT; + HIP_CHECK(hipArrayCreate(&array[i], &desc)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + ARRAY_DESTROY(array[i]); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if ((pavail != avail)) { + HIPASSERT(false); + } + } +} + +/*Thread function*/ +static void ArrayCreateThreadFunc(int gpu) { + ArrayCreate_DiffSizes(gpu); +} + +/* This testcase verifies hipArrayCreate API for small and big chunks data*/ +TEST_CASE("Unit_hipArrayCreate_DiffSizes") { + ArrayCreate_DiffSizes(0); +} + + +/* This testcase verifies the negative scenarios of + * hipArrayCreate API + */ +TEST_CASE("Unit_hipArrayCreate_Negative") { + HIP_ARRAY_DESCRIPTOR desc; + HIP_ARRAY array; + desc.Format = HIP_AD_FORMAT_FLOAT; + desc.NumChannels = 1; + desc.Width = NUM_W; + desc.Height = NUM_H; +#if HT_NVIDIA + SECTION("NullPointer to Array") { + REQUIRE(hipArrayCreate(nullptr, &desc) != hipSuccess); + } + + SECTION("NullPointer to Channel Descriptor") { + REQUIRE(hipArrayCreate(&array, nullptr) != hipSuccess); + } +#endif + SECTION("Width 0 for Array Descriptor") { + desc.Width = 0; + REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess); + } + + SECTION("Invalid NumChannels") { + desc.NumChannels = 3; + REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess); + } +} +/* +This testcase verifies the hipArrayCreate API in multithreaded +scenario by launching threads in parallel on multiple GPUs +and verifies the hipArrayCreate API with small and big chunks data +*/ +TEST_CASE("Unit_hipArrayCreate_MultiThread") { + std::vector threadlist; + int devCnt = 0; + + devCnt = HipTest::getDeviceCount(); + + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < devCnt; i++) { + threadlist.push_back(std::thread(ArrayCreateThreadFunc, i)); + } + + for (auto &t : threadlist) { + t.join(); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + if (pavail != avail) { + WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); + REQUIRE(false); + } +} + diff --git a/catch/unit/memory/hipMalloc3D.cc b/catch/unit/memory/hipMalloc3D.cc new file mode 100644 index 0000000000..2efc9ff7db --- /dev/null +++ b/catch/unit/memory/hipMalloc3D.cc @@ -0,0 +1,133 @@ +/* +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 +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. +*/ + +/* +hipMalloc3D API test scenarios +1. Basic Functionality +2. Negative Scenarios +3. Allocating Small and big chunk data +4. Multithreaded scenario +*/ + +#include +static constexpr auto SMALL_SIZE{4}; +static constexpr auto CHUNK_LOOP{100}; +static constexpr auto BIG_SIZE{100}; +/* +This API verifies hipMalloc3D API by allocating memory in smaller chunks for +CHUNK_LOOP iterations and checks for the memory leaks by get the memory +info before and after the hipMalloc3D API and the difference should +match with the allocated memory +*/ +static void MemoryAlloc3DDiffSizes(int gpu) { + HIPCHECK(hipSetDevice(gpu)); + std::vector array_size; + array_size.push_back(SMALL_SIZE); + array_size.push_back(BIG_SIZE); + for (auto &sizes : array_size) { + size_t width = sizes * sizeof(float); + size_t height{sizes}, depth{sizes}; + hipPitchedPtr devPitchedPtr[CHUNK_LOOP]; + hipExtent extent = make_hipExtent(width, height, depth); + size_t tot, avail, ptot, pavail; + HIPCHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < CHUNK_LOOP; i++) { + HIPCHECK(hipMalloc3D(&devPitchedPtr[i], extent)); + } + for (int i = 0; i < CHUNK_LOOP; i++) { + HIPCHECK(hipFree(devPitchedPtr[i].ptr)); + } + HIPCHECK(hipMemGetInfo(&avail, &tot)); + if ((pavail != avail)) { + HIPASSERT(false); + } + } +} + +static void Malloc3DThreadFunc(int gpu) { + MemoryAlloc3DDiffSizes(gpu); +} + +/* + * This verifies the negative scenarios of hipMalloc3D API + */ +TEST_CASE("Unit_hipMalloc3D_Negative") { + size_t width = SMALL_SIZE * sizeof(char); + size_t height{SMALL_SIZE}, depth{SMALL_SIZE}; + hipPitchedPtr devPitchedPtr; + + SECTION("Passing nullptr to device pitched pointer") { + hipExtent extent = make_hipExtent(width, height, depth); + REQUIRE(hipMalloc3D(nullptr, extent) != hipSuccess); + } + + SECTION("Passing Max values to extent") { + hipExtent extent = make_hipExtent(std::numeric_limits::max(), + std::numeric_limits::max(), + std::numeric_limits::max()); + REQUIRE(hipMalloc3D(&devPitchedPtr, extent) != hipSuccess); + } +} +/* + * This verifies the hipMalloc3D API by + * assigning width,height and depth as 10 + */ +TEST_CASE("Unit_hipMalloc3D_Basic") { + size_t width = SMALL_SIZE * sizeof(char); + size_t height{SMALL_SIZE}, depth{SMALL_SIZE}; + hipPitchedPtr devPitchedPtr; + hipExtent extent = make_hipExtent(width, height, depth); + REQUIRE(hipMalloc3D(&devPitchedPtr, extent) == hipSuccess); +} + +/* +This testcase verifies the hipMalloc3D API by allocating +smaller and big chunk data. +*/ +TEST_CASE("Unit_hipMalloc3D_SmallandBigChunks") { + MemoryAlloc3DDiffSizes(0); +} + +/* +This testcase verifies the hipMalloc3D API in multithreaded +scenario by launching threads in parallel on multiple GPUs +and verifies the hipMalloc3D API with small and big chunks data +*/ +TEST_CASE("Unit_hipMalloc3D_MultiThread") { + std::vector threadlist; + int devCnt = 0; + + devCnt = HipTest::getDeviceCount(); + + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < devCnt; i++) { + threadlist.push_back(std::thread(Malloc3DThreadFunc, i)); + } + + for (auto &t : threadlist) { + t.join(); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + if (pavail != avail) { + WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); + REQUIRE(false); + } +} diff --git a/catch/unit/memory/hipMalloc3DArray.cc b/catch/unit/memory/hipMalloc3DArray.cc new file mode 100644 index 0000000000..fe8ba3f8b0 --- /dev/null +++ b/catch/unit/memory/hipMalloc3DArray.cc @@ -0,0 +1,195 @@ +/* +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 +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. +*/ + +/* +hipMalloc3DArray API test scenarios +1. Basic Functionality +2. Negative Scenarios +3. Allocating Small and big chunk data +4. Multithreaded scenario +*/ + + + +#include + +static constexpr auto ARRAY_SIZE{4}; +static constexpr auto BIG_ARRAY_SIZE{100}; +static constexpr auto ARRAY_LOOP{100}; + +/* + * This API verifies memory allocations for small and + * bigger chunks of data. + * Two scenarios are verified in this API + * 1. SmallArray: Allocates ARRAY_SIZE in a loop and + * releases the memory and verifies the meminfo. + * 2. BigArray: Allocates BIG_ARRAY_SIZE in a loop and + * releases the memory and verifies the meminfo + * + * In both cases, the memory info before allocation and + * after releasing the memory should be the same + * + */ +static void Malloc3DArray_DiffSizes(int gpu) { + HIP_CHECK(hipSetDevice(gpu)); + std::vector array_size; + array_size.push_back(ARRAY_SIZE); + array_size.push_back(BIG_ARRAY_SIZE); + for (auto &size : array_size) { + int width{size}, height{size}, depth{size}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, + 0, 0, hipChannelFormatKindFloat); + hipArray *arr[ARRAY_LOOP]; + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK(hipFreeArray(arr[i])); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if ((pavail != avail)) { + HIPASSERT(false); + } + } +} + +/* Thread Function */ +static void Malloc3DArrayThreadFunc(int gpu) { + Malloc3DArray_DiffSizes(gpu); +} + +/* + * Verifies the negative scenarios of hipMalloc3DArray API + */ +TEST_CASE("Unit_hipMalloc3DArray_Negative") { + constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, + 0, 0, hipChannelFormatKindFloat); + hipArray *arr; +#if HT_NVIDIA + SECTION("NullPointer to Array") { + REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault) != hipSuccess); + } + + SECTION("NullPointer to Channel Descriptor") { + REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width, + height, depth), hipArrayDefault) != hipSuccess); + } +#endif + SECTION("Width 0 in hipExtent") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, + height, width), hipArrayDefault) != hipSuccess); + } + + SECTION("Height 0 in hipExtent") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, + 0, width), hipArrayDefault) != hipSuccess); + } + + SECTION("Invalid Flag") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, + height, depth), 100) != hipSuccess); + } + + SECTION("Width,Height & Depth 0 in hipExtent") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, + 0, 0), hipArrayDefault) != hipSuccess); + } + + SECTION("Max int values to extent") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, + make_hipExtent(std::numeric_limits::max(), + std::numeric_limits::max(), + std::numeric_limits::max()), + hipArrayDefault) != hipSuccess); + } +} +/* + * Verifies the extent validation scenarios + * 1. Passing depth as 0 would create 2D array + * 2. Passing height and depth as 0 would create 1D array + * from hipMalloc3DArray API + */ +TEST_CASE("Unit_hipMalloc3DArray_ExtentValidation") { + constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, + 0, 0, hipChannelFormatKindFloat); + hipArray *arr; + + SECTION("Depth 0 in hipExtent") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, + height, 0), hipArrayDefault) == hipSuccess); + HIP_CHECK(hipFreeArray(arr)); + } + + SECTION("Height & Depth 0 in hipExtent") { + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, + 0, 0), hipArrayDefault) == hipSuccess); + HIP_CHECK(hipFreeArray(arr)); + } +} + +/* + * Verifies hipMalloc3DArray API by passing width,height + * and depth as 10 + */ +TEST_CASE("Unit_hipMalloc3DArray_Basic") { + constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, + 0, 0, hipChannelFormatKindFloat); + hipArray *arr; + + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault) == hipSuccess); + HIP_CHECK(hipFreeArray(arr)); +} + +TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") { + Malloc3DArray_DiffSizes(0); +} +/* +This testcase verifies the hipMalloc3DArray API in multithreaded +scenario by launching threads in parallel on multiple GPUs +and verifies the hipMalloc3DArray API with small and big chunks data +*/ +TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { + std::vector threadlist; + int devCnt = 0; + devCnt = HipTest::getDeviceCount(); + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < devCnt; i++) { + threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i)); + } + + for (auto &t : threadlist) { + t.join(); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + if (pavail != avail) { + WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); + REQUIRE(false); + } +} + diff --git a/catch/unit/memory/hipMallocArray.cc b/catch/unit/memory/hipMallocArray.cc new file mode 100644 index 0000000000..3e521bfcb5 --- /dev/null +++ b/catch/unit/memory/hipMallocArray.cc @@ -0,0 +1,172 @@ +/* +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 +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. +*/ + +/* +hipMallocArray API test scenarios +1. Basic Functionality +2. Negative Scenarios +3. Allocating Small and big chunk data +4. Multithreaded scenario +*/ + + +#include + +static constexpr auto NUM_W{4}; +static constexpr auto BIGNUM_W{100}; +static constexpr auto BIGNUM_H{100}; +static constexpr auto NUM_H{4}; +static constexpr auto ARRAY_LOOP{100}; + +/* + * This API verifies memory allocations for small and + * bigger chunks of data. + * Two scenarios are verified in this API + * 1. NUM_W(small Data): Allocates NUM_W*NUM_H in a loop and + * releases the memory and verifies the meminfo. + * 2. BIGNUM_W(big data): Allocates BIGNUM_W*BIGNUM_H in a loop and + * releases the memory and verifies the meminfo + * + * In both cases, the memory info before allocation and + * after releasing the memory should be the same + * + */ + +static void MallocArray_DiffSizes(int gpu) { + HIP_CHECK(hipSetDevice(gpu)); + std::vector array_size; + array_size.push_back(NUM_W); + array_size.push_back(BIGNUM_W); + for (auto &size : array_size) { + hipArray* A_d[ARRAY_LOOP]; + size_t tot, avail, ptot, pavail; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < ARRAY_LOOP; i++) { + if (size == NUM_W) { + HIP_CHECK(hipMallocArray(&A_d[i], &desc, + NUM_W, NUM_H, + hipArrayDefault)); + } else { + HIP_CHECK(hipMallocArray(&A_d[i], &desc, + BIGNUM_W, BIGNUM_H, + hipArrayDefault)); + } + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK(hipFreeArray(A_d[i])); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if ((pavail != avail)) { + HIPASSERT(false); + } + } +} + +static void MallocArrayThreadFunc(int gpu) { + MallocArray_DiffSizes(gpu); +} + +/* + * This testcase verifies the negative scenarios of + * hipMallocArray API + */ +TEST_CASE("Unit_hipMallocArray_Negative") { + hipArray* A_d; + hipChannelFormatDesc desc = hipCreateChannelDesc(); +#if HT_NVIDIA + SECTION("NullPointer to Array") { + REQUIRE(hipMallocArray(nullptr, &desc, + NUM_W, NUM_H, hipArrayDefault) != hipSuccess); + } + + SECTION("NullPointer to Channel Descriptor") { + REQUIRE(hipMallocArray(&A_d, nullptr, + NUM_W, NUM_H, hipArrayDefault) != hipSuccess); + } +#endif + SECTION("Width 0 in hipMallocArray") { + REQUIRE(hipMallocArray(&A_d, &desc, + 0, NUM_H, hipArrayDefault) != hipSuccess); + } + + SECTION("Height 0 in hipMallocArray") { + REQUIRE(hipMallocArray(&A_d, &desc, + NUM_W, 0, hipArrayDefault) == hipSuccess); + } + + SECTION("Invalid Flag") { + REQUIRE(hipMallocArray(&A_d, &desc, + NUM_W, NUM_H, 100) != hipSuccess); + } + + SECTION("Max int values") { + REQUIRE(hipMallocArray(&A_d, &desc, + std::numeric_limits::max(), + std::numeric_limits::max(), + hipArrayDefault) != hipSuccess); + } +} +/* + * This testcase verifies the basic scenario of + * hipMallocArray API for different datatypes + * of size 10 + */ +TEMPLATE_TEST_CASE("Unit_hipMallocArray_Basic", + "", int, unsigned int, float) { + hipArray* A_d; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + REQUIRE(hipMallocArray(&A_d, &desc, + NUM_W, NUM_H, + hipArrayDefault) == hipSuccess); + HIP_CHECK(hipFreeArray(A_d)); +} + + +TEST_CASE("Unit_hipMallocArray_DiffSizes") { + MallocArray_DiffSizes(0); +} + + +/* +This testcase verifies the hipMallocArray API in multithreaded +scenario by launching threads in parallel on multiple GPUs +and verifies the hipMallocArray API with small and big chunks data +*/ +TEST_CASE("Unit_hipMallocArray_MultiThread") { + std::vector threadlist; + int devCnt = 0; + devCnt = HipTest::getDeviceCount(); + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < devCnt; i++) { + threadlist.push_back(std::thread(MallocArrayThreadFunc, i)); + } + + for (auto &t : threadlist) { + t.join(); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + if (pavail != avail) { + WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); + REQUIRE(false); + } +} + diff --git a/catch/unit/memory/hipMallocPitch.cc b/catch/unit/memory/hipMallocPitch.cc new file mode 100644 index 0000000000..6a25a2f2cb --- /dev/null +++ b/catch/unit/memory/hipMallocPitch.cc @@ -0,0 +1,296 @@ +/* +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 +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 of hipMallocPitch API +1. Negative Scenarios +2. Basic Functionality Scenario +3. Allocate memory using hipMallocPitch API, Launch Kernel validate result. +4. Allocate Memory in small chunks and large chunks and check for possible memory leaks +5. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables. +6. Multithreaded scenario +*/ + +#include +#include +#include + +static constexpr auto SMALLCHUNK_NUMW{4}; +static constexpr auto SMALLCHUNK_NUMH{4}; +static constexpr auto LARGECHUNK_NUMW{1025}; +static constexpr auto LARGECHUNK_NUMH{1000}; +static constexpr auto NUM_W{10}; +static constexpr auto NUM_H{10}; +static constexpr auto COLUMNS{8}; +static constexpr auto ROWS{8}; +static constexpr auto CHUNK_LOOP{100}; + + +template +__global__ void copy_var(T* A, T* B, + size_t ROWS, size_t pitch_A) { + for (uint64_t i = 0; i< ROWS*pitch_A; i= i+pitch_A) { + A[i] = B[i]; + } +} +template +static bool validateResult(T* A, T* B, size_t pitch_A) { + bool testResult = true; + for (uint64_t i=0; i < pitch_A*ROWS; i=i+pitch_A) { + if (A[i] != B[i]) { + testResult = false; + break; + } + } + return testResult; +} +/* + * This API verifies memory allocations for small and + * bigger chunks of data. + * Two scenarios are verified in this API + * 1. SmallChunk: Allocates SMALLCHUNK_NUMW in a loop and + * releases the memory and verifies the meminfo. + * 2. LargeChunk: Allocates LARGECHUNK_NUMW in a loop and + * releases the memory and verifies the meminfo + * + * In both cases, the memory info before allocation and + * after releasing the memory should be the same + * + */ +template +static void MemoryAllocDiffSizes(int gpu) { + HIP_CHECK(hipSetDevice(gpu)); + std::vector array_size; + array_size.push_back(SMALLCHUNK_NUMH); + array_size.push_back(LARGECHUNK_NUMH); + for (auto &sizes : array_size) { + T* A_d[CHUNK_LOOP]; + size_t pitch_A; + size_t width; + if (sizes == SMALLCHUNK_NUMH) { + width = SMALLCHUNK_NUMW * sizeof(T); + } else { + width = LARGECHUNK_NUMW * sizeof(T); + } + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < CHUNK_LOOP; i++) { + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d[i]), + &pitch_A, width, sizes)); + } + for (int i = 0; i < CHUNK_LOOP; i++) { + HIP_CHECK(hipFree(A_d[i])); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if (pavail != avail) { + HIPASSERT(false); + } + } +} + +/*Thread Function */ +static void threadFunc(int gpu) { + MemoryAllocDiffSizes(gpu); +} +/* + * This testcase verifies the negative scenarios of hipMallocPitch API + */ +TEST_CASE("Unit_hipMallocPitch_Negative") { + float* A_d; + size_t pitch_A; + size_t width{NUM_W * sizeof(float)}; +#if HT_NVIDIA + SECTION("NullPtr to Pitched Ptr") { + REQUIRE(hipMallocPitch(nullptr, + &pitch_A, width, NUM_H) != hipSuccess); + } + + SECTION("nullptr to pitch") { + REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), + nullptr, width, NUM_H) != hipSuccess); + } +#endif + SECTION("Width 0 in hipMallocPitch") { + REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, 0, NUM_H) == hipSuccess); + } + + SECTION("Height 0 in hipMallocPitch") { + REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, 0) == hipSuccess); + } + + SECTION("Max int values") { + REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, std::numeric_limits::max(), + std::numeric_limits::max()) != hipSuccess); + } +} +/* + * This testcase verifies the basic scenario of + * hipMallocPitch API for different datatypes + * + */ +TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic", + "[hipMallocPitch]", int, unsigned int, float) { + TestType* A_d; + size_t pitch_A; + size_t width{NUM_W * sizeof(TestType)}; + REQUIRE(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H) == hipSuccess); + HIP_CHECK(hipFree(A_d)); +} + +/* + * This testcase verifies hipMallocPitch API for small + * and big chunks of data. + */ +TEMPLATE_TEST_CASE("Unit_hipMallocPitch_SmallandBigChunks", + "[hipMallocPitch]", int, unsigned int, float) { + MemoryAllocDiffSizes(0); +} + +/* + * This testcase verifies the memory allocated by hipMallocPitch API + * by performing Memcpy2D on the allocated memory. + */ +TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Memcpy2D", "" + , int, float, double) { + HIP_CHECK(hipSetDevice(0)); + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr}, + *B_d{nullptr}; + size_t pitch_A, pitch_B; + size_t width{NUM_W * sizeof(TestType)}; + + // Allocating memory + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, &C_h, NUM_W*NUM_H, false); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), + &pitch_B, width, NUM_H)); + + // Initialize the data + HipTest::setDefaultData(NUM_W*NUM_H, A_h, B_h, C_h); + + // Host to Device + HIP_CHECK(hipMemcpy2D(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), + COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice)); + + // Performs D2D on same GPU device + HIP_CHECK(hipMemcpy2D(B_d, pitch_B, A_d, + pitch_A, COLUMNS*sizeof(TestType), + ROWS, hipMemcpyDeviceToDevice)); + + // hipMemcpy2D Device to Host + HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyDeviceToHost)); + + // Validating the result + REQUIRE(HipTest::checkArray(A_h, B_h, COLUMNS, ROWS) == true); + + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +} + + + +/* +This testcase verifies the hipMallocPitch API in multithreaded +scenario by launching threads in parallel on multiple GPUs +and verifies the hipMallocPitch API with small and big chunks data +*/ + +TEST_CASE("Unit_hipMallocPitch_MultiThread", "") { + std::vector threadlist; + int devCnt = 0; + + devCnt = HipTest::getDeviceCount(); + + size_t tot, avail, ptot, pavail; + HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + for (int i = 0; i < devCnt; i++) { + threadlist.push_back(std::thread(threadFunc, i)); + } + + for (auto &t : threadlist) { + t.join(); + } + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + + if (pavail != avail) { + WARN("Memory leak of hipMallocPitch API in multithreaded scenario"); + REQUIRE(false); + } +} + +/* + * This testcase verifies hipMallocPitch API by + * 1. Allocating Memory using hipMallocPitch API + * 2. Launching the kernel and copying the data from the allocated kernel + * variable to another kernel variable. + * 3. Validating the result + */ +TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", "" + , int, float, double) { + HIP_CHECK(hipSetDevice(0)); + TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr}, + *B_d{nullptr}; + size_t pitch_A, pitch_B; + size_t width{NUM_W * sizeof(TestType)}; + + // Allocating memory + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, &C_h, NUM_W*NUM_H, false); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), + &pitch_B, width, NUM_H)); + + // Host to Device + HIP_CHECK(hipMemcpy2D(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), + COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice)); + + + hipLaunchKernelGGL(copy_var, dim3(1), dim3(1), + 0, 0, static_cast(A_d), + static_cast(B_d), ROWS, pitch_A); + + + // hipMemcpy2D Device to Host + HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyDeviceToHost)); + + // Validating the result + validateResult(A_h, B_h, pitch_A); + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +} + +