From d58953d4e99b13471eefd655e04f3d75fcc83820 Mon Sep 17 00:00:00 2001 From: Finlay Date: Mon, 11 Jul 2022 07:49:05 +0100 Subject: [PATCH] EXSWCPHIPT-112 - Unit test for hipMalloc3DArray for default and surface arrays (#2713) [ROCm/hip commit: c0deb17bbcd5bfdbbba0c7d456625f28e6e45015] --- .../tests/catch/unit/memory/hipArrayCommon.hh | 36 +++ .../tests/catch/unit/memory/hipArrayCreate.cc | 89 +++--- .../catch/unit/memory/hipMalloc3DArray.cc | 266 ++++++++++++------ .../tests/catch/unit/memory/hipMallocArray.cc | 95 +++---- 4 files changed, 302 insertions(+), 184 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipArrayCommon.hh b/projects/hip/tests/catch/unit/memory/hipArrayCommon.hh index 1c6100a6e9..c9823806c5 100644 --- a/projects/hip/tests/catch/unit/memory/hipArrayCommon.hh +++ b/projects/hip/tests/catch/unit/memory/hipArrayCommon.hh @@ -122,3 +122,39 @@ inline size_t getFreeMem() { HIP_CHECK(hipMemGetInfo(&free, &total)); return free; } + +struct Sizes { + int max1D; + std::array max2D; + std::array max3D; + + Sizes(unsigned int flag) { + int device; + HIP_CHECK(hipGetDevice(&device)); + switch (flag) { + case hipArrayDefault: { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + max1D = prop.maxTexture1D; + max2D = {prop.maxTexture2D[0], prop.maxTexture2D[1]}; + max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]}; + return; + } + case hipArraySurfaceLoadStore: { + int value; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device)); + max1D = value; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface2D, device)); + max2D = {value, value}; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface3D, device)); + max3D = {value, value, value}; + return; + } + default: { + INFO("Array flag not supported"); + REQUIRE(false); + return; + } + } + } +}; diff --git a/projects/hip/tests/catch/unit/memory/hipArrayCreate.cc b/projects/hip/tests/catch/unit/memory/hipArrayCreate.cc index 9ef637821e..d55bdae255 100644 --- a/projects/hip/tests/catch/unit/memory/hipArrayCreate.cc +++ b/projects/hip/tests/catch/unit/memory/hipArrayCreate.cc @@ -30,10 +30,10 @@ hipArrayCreate API test scenarios #include "hipArrayCommon.hh" #include "DriverContext.hh" -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 size_t NUM_W{4}; +static constexpr size_t BIGNUM_W{100}; +static constexpr size_t NUM_H{4}; +static constexpr size_t BIGNUM_H{100}; static constexpr auto ARRAY_LOOP{100}; /* @@ -51,31 +51,34 @@ static constexpr auto ARRAY_LOOP{100}; */ static void ArrayCreate_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; - for (auto& size : array_size) { - std::array array; - const size_t pavail = getFreeMem(); - HIP_ARRAY_DESCRIPTOR desc; - desc.NumChannels = 1; - desc.Width = std::get<0>(size); - desc.Height = std::get<1>(size); - desc.Format = HIP_AD_FORMAT_FLOAT; - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipArrayCreate(&array[i], &desc)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipArrayDestroy(array[i])); - } - const size_t avail = getFreeMem(); - if (pavail != avail) { - HIPASSERT(false); - } + HIP_CHECK_THREAD(hipSetDevice(gpu)); + std::pair size = + GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); + std::array array; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + HIP_ARRAY_DESCRIPTOR desc; + desc.NumChannels = 1; + desc.Width = std::get<0>(size); + desc.Height = std::get<1>(size); + desc.Format = HIP_AD_FORMAT_FLOAT; + + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc)); } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayDestroy(array[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } /* This testcase verifies hipArrayCreate API for small and big chunks data*/ -TEST_CASE("Unit_hipArrayCreate_DiffSizes") { ArrayCreate_DiffSizes(0); } +TEST_CASE("Unit_hipArrayCreate_DiffSizes") { + ArrayCreate_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); +} /* This testcase verifies the hipArrayCreate API in multithreaded @@ -90,13 +93,13 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { const size_t pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(ArrayCreate_DiffSizes, i)); } for (auto& t : threadlist) { t.join(); } + HIP_CHECK_THREAD_FINALIZE(); const size_t avail = getFreeMem(); if (pavail != avail) { @@ -282,28 +285,26 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort desc.Format = vec_info::format; desc.NumChannels = vec_info::size; - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t prop; - HIP_CHECK(hipGetDeviceProperties(&prop, device)); + const Sizes sizes(hipArrayDefault); + const size_t s = 64; hiparray array{}; SECTION("Happy") { SECTION("1D - Max") { - desc.Width = prop.maxTexture1D; + desc.Width = sizes.max1D; desc.Height = 0; } SECTION("2D - Max Width") { - desc.Width = prop.maxTexture2D[0]; - desc.Height = 64; + desc.Width = sizes.max2D[0]; + desc.Height = s; } SECTION("2D - Max Height") { - desc.Width = 64; - desc.Height = prop.maxTexture2D[1]; + desc.Width = s; + desc.Height = sizes.max2D[1]; } SECTION("2D - Max Width and Height") { - desc.Width = prop.maxTexture2D[0]; - desc.Height = prop.maxTexture2D[1]; + desc.Width = sizes.max2D[0]; + desc.Height = sizes.max2D[1]; } auto maxArrayCreateError = hipArrayCreate(&array, &desc); // this can try to alloc many GB of memory, so out of memory is acceptable @@ -314,20 +315,20 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort } SECTION("Negative") { SECTION("1D - More Than Max") { - desc.Width = prop.maxTexture1D + 1; + desc.Width = sizes.max1D + 1; desc.Height = 0; } SECTION("2D - More Than Max Width") { - desc.Width = prop.maxTexture2D[0] + 1; - desc.Height = 64; + desc.Width = sizes.max2D[0] + 1; + desc.Height = s; } SECTION("2D - More Than Max Height") { - desc.Width = 64; - desc.Height = prop.maxTexture2D[1] + 1; + desc.Width = s; + desc.Height = sizes.max2D[1] + 1; } SECTION("2D - More Than Max Width and Height") { - desc.Width = prop.maxTexture2D[0] + 1; - desc.Height = prop.maxTexture2D[1] + 1; + desc.Width = sizes.max2D[0] + 1; + desc.Height = sizes.max2D[1] + 1; } HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); } diff --git a/projects/hip/tests/catch/unit/memory/hipMalloc3DArray.cc b/projects/hip/tests/catch/unit/memory/hipMalloc3DArray.cc index fe8ba3f8b0..253913345e 100644 --- a/projects/hip/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/projects/hip/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -25,9 +25,9 @@ hipMalloc3DArray API test scenarios 4. Multithreaded scenario */ - - +#include #include +#include "hipArrayCommon.hh" static constexpr auto ARRAY_SIZE{4}; static constexpr auto BIG_ARRAY_SIZE{100}; @@ -47,34 +47,24 @@ static constexpr auto ARRAY_LOOP{100}; * */ 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); - } - } -} + HIP_CHECK_THREAD(hipSetDevice(gpu)); + const int size = GENERATE(ARRAY_SIZE, BIG_ARRAY_SIZE); + int width{size}, height{size}, depth{size}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + std::array arr; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); -/* Thread Function */ -static void Malloc3DArrayThreadFunc(int gpu) { - Malloc3DArray_DiffSizes(gpu); + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth), + hipArrayDefault)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(arr[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } /* @@ -82,91 +72,53 @@ static void Malloc3DArrayThreadFunc(int gpu) { */ 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; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray* arr; #if HT_NVIDIA SECTION("NullPointer to Array") { - REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault) != hipSuccess); + 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); + 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); + 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); + 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); + 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); + 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); + 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); + HIP_CHECK_THREAD_FINALIZE(); } + /* This testcase verifies the hipMalloc3DArray API in multithreaded scenario by launching threads in parallel on multiple GPUs @@ -176,16 +128,16 @@ 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)); + const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i)); + threadlist.push_back(std::thread(Malloc3DArray_DiffSizes, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const auto avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); @@ -193,3 +145,131 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { } } +void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc, + const hipExtent& expected_extent, const unsigned int expected_flags) { +// hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) +#if HT_AMD + std::ignore = array; + std::ignore = expected_desc; + std::ignore = expected_extent; + std::ignore = expected_flags; +#else + cudaChannelFormatDesc queried_desc; + cudaExtent queried_extent; + unsigned int queried_flags; + + cudaArrayGetInfo(&queried_desc, &queried_extent, &queried_flags, array); + + REQUIRE(expected_desc.x == queried_desc.x); + REQUIRE(expected_desc.y == queried_desc.y); + REQUIRE(expected_desc.z == queried_desc.z); + REQUIRE(expected_desc.f == queried_desc.f); + + REQUIRE(expected_extent.width == queried_extent.width); + REQUIRE(expected_extent.height == queried_extent.height); + REQUIRE(expected_extent.depth == queried_extent.depth); + + REQUIRE(expected_flags == queried_flags); +#endif +} + +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, + float2, float4) { + hipArray_t array; + const auto desc = hipCreateChannelDesc(); +#if HT_AMD + const unsigned int flags = hipArrayDefault; +#else + const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); +#endif + constexpr size_t size = 64; + hipExtent extent; + + SECTION("1D Array") { + extent = make_hipExtent(size, 0, 0); + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + } + SECTION("2D Array") { + extent = make_hipExtent(size, size, 0); + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + } + SECTION("3D Array") { + extent = make_hipExtent(size, size, size); + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + } + + checkArrayIsExpected(array, desc, extent, flags); + + HIP_CHECK(hipFreeArray(array)); +} + +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, ushort2, + unsigned char, float, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97"); + return; +#endif + + hipArray_t array; + const hipChannelFormatDesc desc = hipCreateChannelDesc(); +#if HT_AMD + const unsigned int flag = hipArrayDefault; +#else + const unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); +#endif + if (flag == hipArraySurfaceLoadStore) { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); + return; + } + CAPTURE(flag); + const Sizes sizes(flag); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { + // stored in a vector so some values can be ifdef'd out + std::vector extentsToTest{ + make_hipExtent(sizes.max1D, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0], s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1], 0), // 2D max height + make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0), // 2D max + make_hipExtent(sizes.max3D[0], s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1], s), // 3D max height + make_hipExtent(s, s, sizes.max3D[2]), // 3D max depth + make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]), // 3D max height and depth + make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]), // 3D max width and depth + make_hipExtent(sizes.max3D[0], sizes.max3D[1], s), // 3D max width and height + make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2]) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + CAPTURE(extent.width, extent.height, extent.depth); + auto maxArrayCreateError = hipMalloc3DArray(&array, &desc, extent, flag); + // this can try to alloc many GB of memory, so out of memory is acceptable + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + checkArrayIsExpected(array, desc, extent, flag); + HIP_CHECK(hipFreeArray(array)); + } + SECTION("Negative") { + std::vector extentsToTest { + make_hipExtent(sizes.max1D + 1, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0] + 1, s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1] + 1, 0), // 2D max height + make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0), // 2D max + make_hipExtent(sizes.max3D[0] + 1, s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1] + 1, s), // 3D max height +#if !HT_NVIDIA // leads to hipSuccess on NVIDIA + make_hipExtent(s, s, sizes.max3D[2] + 1), // 3D max depth +#endif + make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1), // 3D max height and depth + make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1), // 3D max width and depth + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s), // 3D max width and height + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + CAPTURE(extent.width, extent.height, extent.depth); + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flag), hipErrorInvalidValue); + } +} diff --git a/projects/hip/tests/catch/unit/memory/hipMallocArray.cc b/projects/hip/tests/catch/unit/memory/hipMallocArray.cc index 8859aadb45..820e917313 100644 --- a/projects/hip/tests/catch/unit/memory/hipMallocArray.cc +++ b/projects/hip/tests/catch/unit/memory/hipMallocArray.cc @@ -30,11 +30,11 @@ hipMallocArray API test scenarios #include #include "hipArrayCommon.hh" -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}; +static constexpr size_t NUM_W{4}; +static constexpr size_t NUM_H{4}; +static constexpr size_t BIGNUM_W{100}; +static constexpr size_t BIGNUM_H{100}; +static constexpr int ARRAY_LOOP{100}; /* * This API verifies memory allocations for small and @@ -50,28 +50,30 @@ static constexpr auto ARRAY_LOOP{100}; * */ static void MallocArray_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; - for (auto& size : array_size) { - std::array A_d; - size_t tot, avail, ptot, pavail; - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK( - hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), 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); - } + HIP_CHECK_THREAD(hipSetDevice(gpu)); + std::pair size = + GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array A_d; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD( + hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(A_d[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } -TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); } +TEST_CASE("Unit_hipMallocArray_DiffSizes") { + MallocArray_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); +} /* This testcase verifies the hipMallocArray API in multithreaded @@ -82,18 +84,17 @@ 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)); + const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const auto avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); REQUIRE(false); @@ -450,32 +451,32 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho // EXSWCPHIPT-71 - no equivalent value for maxSurface and maxTexture2DGather. TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ushort, short2, char, char4, float2, float4) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t prop; - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - size_t width, height; hipArray_t array{}; hipChannelFormatDesc desc = hipCreateChannelDesc(); const unsigned int flag = hipArrayDefault; + const Sizes sizes(flag); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { SECTION("1D - Max") { - width = prop.maxTexture1D; + width = sizes.max1D; height = 0; } SECTION("2D - Max Width") { - width = prop.maxTexture2D[0]; - height = 64; + width = sizes.max2D[0]; + height = s; } SECTION("2D - Max Height") { - width = 64; - height = prop.maxTexture2D[1]; + width = s; + height = sizes.max2D[1]; } SECTION("2D - Max Width and Height") { - width = prop.maxTexture2D[0]; - height = prop.maxTexture2D[1]; + width = sizes.max2D[0]; + height = sizes.max2D[1]; } auto maxArrayCreateError = hipMallocArray(&array, &desc, width, height, flag); // this can try to alloc many GB of memory, so out of memory is acceptable @@ -485,20 +486,20 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush } SECTION("Negative") { SECTION("1D - More Than Max") { - width = prop.maxTexture1D + 1; + width = sizes.max1D + 1; height = 0; } SECTION("2D - More Than Max Width") { - width = prop.maxTexture2D[0] + 1; - height = 64; + width = sizes.max2D[0] + 1; + height = s; } SECTION("2D - More Than Max Height") { - width = 64; - height = prop.maxTexture2D[1] + 1; + width = s; + height = sizes.max2D[1] + 1; } SECTION("2D - More Than Max Width and Height") { - width = prop.maxTexture2D[0] + 1; - height = prop.maxTexture2D[1] + 1; + width = sizes.max2D[0] + 1; + height = sizes.max2D[1] + 1; } HIP_CHECK_ERROR(hipMallocArray(&array, &desc, width, height, flag), hipErrorInvalidValue); }