diff --git a/projects/hip-tests/catch/unit/memory/hipMallocArray.cc b/projects/hip-tests/catch/unit/memory/hipMallocArray.cc index a73566b40d..b87c9209b8 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocArray.cc +++ b/projects/hip-tests/catch/unit/memory/hipMallocArray.cc @@ -26,6 +26,7 @@ hipMallocArray API test scenarios */ #include +#include static constexpr auto NUM_W{4}; static constexpr auto BIGNUM_W{100}; @@ -46,23 +47,17 @@ static constexpr auto ARRAY_LOOP{100}; * 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); + std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; for (auto& size : array_size) { - hipArray* A_d[ARRAY_LOOP]; + 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++) { - 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)); - } + 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])); @@ -74,42 +69,8 @@ static void MallocArray_DiffSizes(int 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); - } -} - - 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 @@ -122,6 +83,7 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { size_t tot, avail, ptot, pavail; HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < devCnt; i++) { + // TODO the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } @@ -236,6 +198,19 @@ template void checkDataIsAscending(const std::vector& hostData) REQUIRE(allMatch); } +const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { + switch (formatKind) { + case hipChannelFormatKindFloat: + return "float"; + case hipChannelFormatKindSigned: + return "signed"; + case hipChannelFormatKindUnsigned: + return "unsigned"; + default: + return "error"; + } +} + // Tests ///////////////////////////////////////// // Test the default array by generating a texture from it then reading from that texture. @@ -363,6 +338,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho char4, float, float2, float4) { #if HT_AMD HipTest::HIP_SKIP_TEST("EXSWCPHIPT-62"); + return; #endif hipChannelFormatDesc desc = hipCreateChannelDesc(); @@ -381,7 +357,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault)); testArrayAsTexture(arrayPtr, width, height); } -#if HT_NVIDIA // surfaces and texture gather not supported on AMD +#if HT_NVIDIA // surfaces not supported on AMD SECTION("hipArraySurfaceLoadStore") { INFO("flag is hipArraySurfaceLoadStore"); INFO("height: " << height); @@ -399,3 +375,308 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho HIP_CHECK(hipFreeArray(arrayPtr)); } + +// Arrays can be up to the size of maxTexture* but no bigger +// 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; + + SECTION("Happy") { + SECTION("1D - Max") { + width = prop.maxTexture1D; + height = 0; + } + SECTION("2D - Max Width") { + width = prop.maxTexture2D[0]; + height = 64; + } + SECTION("2D - Max Height") { + width = 64; + height = prop.maxTexture2D[1]; + } + SECTION("2D - Max Width and Height") { + width = prop.maxTexture2D[0]; + height = prop.maxTexture2D[1]; + } + auto maxArrayCreateError = hipMallocArray(&array, &desc, width, height, flag); + // this can try to alloc many GB of memory, so out of memory is fair + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + HIP_CHECK(hipFreeArray(array)); + } + SECTION("Negative") { + SECTION("1D - More Than Max") { + width = prop.maxTexture1D + 1; + height = 0; + } + SECTION("2D - More Than Max Width") { + width = prop.maxTexture2D[0] + 1; + height = 64; + } + SECTION("2D - More Than Max Height") { + width = 64; + height = prop.maxTexture2D[1] + 1; + } + SECTION("2D - More Than Max Width and Height") { + width = prop.maxTexture2D[0] + 1; + height = prop.maxTexture2D[1] + 1; + } + HIP_CHECK_ERROR(hipMallocArray(&array, &desc, width, height, flag), hipErrorInvalidValue); + } +} + + +// Arrays with channels of different size are not allowed. +TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + return; +#endif + const int bitsX = GENERATE(8, 16, 32); + const int bitsY = GENERATE(8, 16, 32); + const int bitsZ = GENERATE(8, 16, 32); + const int bitsW = GENERATE(8, 16, 32); + if (bitsX == bitsY && bitsY == bitsZ && bitsZ == bitsW) return; // skip when they are equal + + const hipChannelFormatKind channelFormat = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + + if (channelFormat == hipChannelFormatKindFloat && + (bitsX == 8 || bitsY == 8 || bitsZ == 8 || bitsW == 8)) + return; // 8 bit floats aren't allowed + + + hipChannelFormatDesc desc = hipCreateChannelDesc(bitsX, bitsY, bitsZ, bitsW, channelFormat); + REQUIRE(desc.x == bitsX); + REQUIRE(desc.y == bitsY); + REQUIRE(desc.z == bitsZ); + REQUIRE(desc.w == bitsW); + + hipArray_t arrayPtr{}; + size_t width = 1024; + size_t height = 1024; + + INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX + << ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW); + +#if HT_AMD + unsigned int flag = hipArrayDefault; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, width, height, flag), hipErrorInvalidValue); +#else + unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, width, height, flag), hipErrorUnknown); +#endif +} + +// Zero-width array is not supported +TEST_CASE("Unit_hipMallocArray_Negative_ZeroWidth") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + // pointer to the array in device memory + hipArray_t arrayPtr; + + size_t width = 0; + size_t height = GENERATE(0, 32); + + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault), + hipErrorInvalidValue); +} + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-45"); + return; +#endif + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + HIP_CHECK_ERROR(hipMallocArray(nullptr, &desc, 1024, 0, hipArrayDefault), hipErrorInvalidValue); +} + +// Providing the desc pointer as nullptr should return an error +TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-83"); + return; +#endif + hipArray_t arrayPtr; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, nullptr, 1024, 0, hipArrayDefault), + hipErrorInvalidValue); +} + +// Inappropriate but related flags should still return an error +TEST_CASE("Unit_hipMallocArray_Negative_BadFlags") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-72"); + return; +#endif + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + hipArray_t arrayPtr; + SECTION("Flags that dont work with 1D") { +#if HT_AMD + // * cudaArrayLayered 0x01 - 1 + // * cudaArrayCubemap 0x04 - 4 + unsigned int flag = + GENERATE(hipArrayLayered, hipArrayCubemap, hipArrayLayered | hipArrayCubemap); +#else + // * cudaArrayTextureGather 0x08 - 8 (2D only) + unsigned int flag = GENERATE(hipArrayTextureGather, hipArrayLayered, hipArrayCubemap, + hipArrayLayered | hipArrayCubemap); +#endif + INFO("Using flag " << flag); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 0, flag), hipErrorInvalidValue); + } + SECTION("Flags that dont work with 2D") { + unsigned int flag = GENERATE(hipArrayCubemap, hipArrayLayered | hipArrayCubemap); + INFO("Using flag " << flag); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); + } +} + +// 8-bit float channels are not supported +TEMPLATE_TEST_CASE("Unit_hipMallocArray_Negative_8bitFloat", "", float, float2, float4) { + hipChannelFormatDesc desc = GENERATE(hipCreateChannelDesc(8, 0, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 8, 8, hipChannelFormatKindFloat)); + + // pointer to the array in device memory + hipArray_t arrayPtr; + +#if HT_AMD + unsigned int flags = hipArrayDefault; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flags), hipErrorInvalidValue); +#else + unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flags), hipErrorUnknown); +#endif +} + +// Only 8, 16, and 32 bit channels are supported +TEST_CASE("Unit_hipMallocArray_Negative_BadNumberOfBits") { + const int badBits = GENERATE(-1, 0, 10, 100); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + hipChannelFormatDesc desc = hipCreateChannelDesc(badBits, badBits, badBits, badBits, formatKind); + + REQUIRE(desc.x == badBits); + REQUIRE(desc.y == badBits); + REQUIRE(desc.z == badBits); + REQUIRE(desc.w == badBits); + + // pointer to the array in device memory + hipArray_t arrayPtr; + + INFO("Number of bits: " << badBits); +#if HT_AMD + unsigned int flag = hipArrayDefault; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); +#else + unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); + INFO("flag: " << flag); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorUnknown); +#endif +} + +// creating elements with 3 channels is not supported. +TEST_CASE("Unit_hipMallocArray_Negative_3ChannelElement") { + const int bits = GENERATE(8, 16, 32); + hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + if (bits == 8 && formatKind == hipChannelFormatKindFloat) return; + + hipChannelFormatDesc desc = hipCreateChannelDesc(bits, bits, bits, 0, formatKind); + + REQUIRE(desc.x == bits); + REQUIRE(desc.y == bits); + REQUIRE(desc.z == bits); + REQUIRE(desc.w == 0); + + // pointer to the array in device memory + hipArray_t arrayPtr; + +#if HT_AMD + unsigned int flag = hipArrayDefault; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); +#else + unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorUnknown); +#endif +} + +// The bit channel description should not allow any channels after a zero channel +TEST_CASE("Unit_hipMallocArray_Negative_ChannelAfterZeroChannel") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + return; +#endif + const int bits = GENERATE(8, 16, 32); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + if (bits == 8 && formatKind == hipChannelFormatKindFloat) return; + + hipChannelFormatDesc desc = GENERATE_COPY(hipCreateChannelDesc(0, bits, bits, 0, formatKind), + hipCreateChannelDesc(0, bits, bits, bits, formatKind), + hipCreateChannelDesc(bits, 0, bits, 0, formatKind), + hipCreateChannelDesc(bits, bits, 0, bits, formatKind), + hipCreateChannelDesc(0, 0, bits, 0, formatKind), + hipCreateChannelDesc(0, 0, bits, bits, formatKind)); + + INFO("kind: " << channelFormatString(formatKind)); + INFO("x: " << desc.x << ", y: " << desc.y << ", z: " << desc.z << ", w: " << desc.w); + + hipArray_t arrayPtr; +#if HT_AMD + unsigned int flag = hipArrayDefault; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); +#else + unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorUnknown); +#endif +} + +// The channel format should be one of the defined formats +TEST_CASE("Unit_hipMallocArray_Negative_InvalidChannelFormat") { + const int bits = 32; + hipChannelFormatKind formatKind = static_cast(0xFF); + hipChannelFormatDesc desc = hipCreateChannelDesc(bits, bits, bits, bits, formatKind); + + REQUIRE(desc.f != hipChannelFormatKindFloat); + REQUIRE(desc.f != hipChannelFormatKindUnsigned); + REQUIRE(desc.f != hipChannelFormatKindSigned); + + hipArray_t arrayPtr; + + CAPTURE(formatKind); + +#if HT_AMD + unsigned int flag = hipArrayDefault; + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); +#else + unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorUnknown); +#endif +} + + +// hipMallocArray should handle the max numeric value gracefully. +TEST_CASE("Unit_hipMallocArray_Negative_NumericLimit") { + hipArray_t arrayPtr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + size_t size = std::numeric_limits::max(); +#if HT_AMD + unsigned int flag = hipArrayDefault; +#else + unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); +#endif + HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, size, size, flag), hipErrorInvalidValue); +}