From daceb2e1a1608e25f8bec631cbf88c40fbded27f Mon Sep 17 00:00:00 2001 From: Finlay Date: Wed, 13 Jul 2022 06:16:26 +0100 Subject: [PATCH] EXSWCPHIPT-134 - Negative tests for hipMalloc3DArray and hipArray3DCreate (#2784) --- tests/catch/unit/memory/hipArray3DCreate.cc | 157 +++++++++++- tests/catch/unit/memory/hipArrayCommon.hh | 43 ++++ tests/catch/unit/memory/hipArrayCreate.cc | 42 +--- tests/catch/unit/memory/hipMalloc3DArray.cc | 262 +++++++++++++++----- tests/catch/unit/memory/hipMallocArray.cc | 32 +-- 5 files changed, 413 insertions(+), 123 deletions(-) diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc index 47e70e66c5..faa3822b9b 100644 --- a/tests/catch/unit/memory/hipArray3DCreate.cc +++ b/tests/catch/unit/memory/hipArray3DCreate.cc @@ -17,8 +17,10 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include #include "DriverContext.hh" #include "hipArrayCommon.hh" +#include "hip_test_common.hh" namespace { void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) { @@ -38,6 +40,11 @@ void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& ex REQUIRE(queried_desc.Flags == expected_desc.Flags); #endif } + +void testInvalidDescription(HIP_ARRAY3D_DESCRIPTOR desc) { + hiparray array; + HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); +} } // namespace TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float, @@ -45,7 +52,6 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, using vec_info = vector_info; DriverContext ctx; - hiparray array; HIP_ARRAY3D_DESCRIPTOR desc{}; desc.Format = vec_info::format; desc.NumChannels = vec_info::size; @@ -70,6 +76,7 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, CAPTURE(desc.Width, desc.Height, desc.Depth); + hiparray array; HIP_CHECK(hipArray3DCreate(&array, &desc)); checkArrayIsExpected(array, desc); HIP_CHECK(hipArrayDestroy(array)); @@ -164,3 +171,151 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, us HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); } } + +#if HT_NVIDIA +constexpr std::array validFlags{ + 0, + hipArraySurfaceLoadStore, + hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayCubemap, + hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayTextureGather}; +#else +constexpr std::array validFlags{ + 0, hipArrayCubemap, hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore}; +#endif + +constexpr HIP_ARRAY3D_DESCRIPTOR defaultDescriptor(unsigned int flags, size_t size) { + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = HIP_AD_FORMAT_FLOAT; + desc.NumChannels = 4; + desc.Flags = flags; + desc.Width = size; + desc.Height = size; + desc.Depth = size; + +#if HT_NVIDIA + if (flags == CUDA_ARRAY3D_TEXTURE_GATHER) { + desc.Depth = 0; + } +#endif + return desc; +} + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipArray3DCreate_Negative_NullArrayPtr") { + auto desc = defaultDescriptor(0, 64); + + DriverContext ctx; + HIP_CHECK_ERROR(hipArray3DCreate(nullptr, &desc), hipErrorInvalidValue); +} + +// Providing the description pointer as nullptr should return an error +TEST_CASE("Unit_hipArray3DCreate_Negative_NullDescPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130"); + return; +#endif + + DriverContext ctx; + hiparray array; + HIP_CHECK_ERROR(hipArray3DCreate(&array, nullptr), hipErrorInvalidValue); +} + + +// Zero width arrays are not allowed +TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroWidth") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + desc.Width = 0; + CAPTURE(desc.Flags); + + testInvalidDescription(desc); +} + +// Zero height arrays are only allowed for 1D arrays and layered arrays +TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroHeight") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); +#if HT_NVIDIA + std::array exceptions{CUDA_ARRAY3D_LAYERED, + CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST}; +#else + std::array exceptions{}; +#endif + desc.Height = 0; + + if (std::find(std::begin(exceptions), std::end(exceptions), desc.Flags) == std::end(exceptions)) { + // flag is not in list of exceptions + testInvalidDescription(desc); + } +} + +// Arrays must be created with a valid data format +TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFormat") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + + desc.Format = static_cast(0xDEADBEEF); + REQUIRE(std::find(std::begin(driverFormats), std::end(driverFormats), desc.Format) == + std::end(driverFormats)); + + testInvalidDescription(desc); +} + +// An array must have either 1,2, or 4 channels +TEST_CASE("Unit_hipArray3DCreate_Negative_NumChannels") { + DriverContext ctx; + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + desc.NumChannels = GENERATE(0, 3, 5); + + testInvalidDescription(desc); +} + +// Using invalid flags should result in an error +TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFlags") { + DriverContext ctx; + + // FIXME: use the same flags for both tests when the values exist for hip +#if HT_NVIDIA + unsigned int flags = + GENERATE(0xDEADBEEF, CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST, + CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_CUBEMAP, + CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_CUBEMAP); +#else + unsigned int flags = 0xDEADBEEF; +#endif + + CAPTURE(flags); + + auto desc = defaultDescriptor(flags, 6); + + + REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), desc.Flags) == + std::end(validFlags)); + + testInvalidDescription(desc); +} + + +// hipArray3DCreate should handle the max numeric value gracefully. +TEST_CASE("Unit_hipArray3DCreate_Negative_NumericLimit") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, std::numeric_limits::max()); + + testInvalidDescription(desc); +} diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh index b51022eb56..8502ac218d 100644 --- a/tests/catch/unit/memory/hipArrayCommon.hh +++ b/tests/catch/unit/memory/hipArrayCommon.hh @@ -166,3 +166,46 @@ struct Sizes { } } }; + +inline const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { + switch (formatKind) { + case hipChannelFormatKindFloat: + return "float"; + case hipChannelFormatKindSigned: + return "signed"; + case hipChannelFormatKindUnsigned: + return "unsigned"; + default: + return "error"; + } +} + +// All the possible formats for channel data in an array. +static const std::vector driverFormats{ + HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32, + HIP_AD_FORMAT_SIGNED_INT8, HIP_AD_FORMAT_SIGNED_INT16, HIP_AD_FORMAT_SIGNED_INT32, + HIP_AD_FORMAT_HALF, HIP_AD_FORMAT_FLOAT}; + +// Helpful for printing errors +inline const char* formatToString(hipArray_Format f) { + switch (f) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + return "Unsigned Int 8"; + case HIP_AD_FORMAT_UNSIGNED_INT16: + return "Unsigned Int 16"; + case HIP_AD_FORMAT_UNSIGNED_INT32: + return "Unsigned Int 32"; + case HIP_AD_FORMAT_SIGNED_INT8: + return "Signed Int 8"; + case HIP_AD_FORMAT_SIGNED_INT16: + return "Signed Int 16"; + case HIP_AD_FORMAT_SIGNED_INT32: + return "Signed Int 32"; + case HIP_AD_FORMAT_HALF: + return "Float 16"; + case HIP_AD_FORMAT_FLOAT: + return "Float 32"; + default: + return "not found"; + } +} diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc index d55bdae255..836c0d1c45 100644 --- a/tests/catch/unit/memory/hipArrayCreate.cc +++ b/tests/catch/unit/memory/hipArrayCreate.cc @@ -109,36 +109,6 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { } -// All the possible formats for channel data in an array. -static const std::vector formats{ - HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32, - HIP_AD_FORMAT_SIGNED_INT8, HIP_AD_FORMAT_SIGNED_INT16, HIP_AD_FORMAT_SIGNED_INT32, - HIP_AD_FORMAT_HALF, HIP_AD_FORMAT_FLOAT}; - -// Helpful for printing errors -const char* formatToString(hipArray_Format f) { - switch (f) { - case HIP_AD_FORMAT_UNSIGNED_INT8: - return "Unsigned Int 8"; - case HIP_AD_FORMAT_UNSIGNED_INT16: - return "Unsigned Int 16"; - case HIP_AD_FORMAT_UNSIGNED_INT32: - return "Unsigned Int 32"; - case HIP_AD_FORMAT_SIGNED_INT8: - return "Signed Int 8"; - case HIP_AD_FORMAT_SIGNED_INT16: - return "Signed Int 16"; - case HIP_AD_FORMAT_SIGNED_INT32: - return "Signed Int 32"; - case HIP_AD_FORMAT_HALF: - return "Float 16"; - case HIP_AD_FORMAT_FLOAT: - return "Float 32"; - default: - return "not found"; - } -} - // Tests ///////////////////////////////////////// #if HT_AMD @@ -338,7 +308,7 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; - desc.Format = formats[0]; + desc.Format = driverFormats[0]; desc.NumChannels = 4; desc.Width = 0; desc.Height = GENERATE(0, 1024); @@ -351,13 +321,13 @@ TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { // HipArrayCreate will return an error when nullptr is used as the array argument TEST_CASE("Unit_hipArrayCreate_Nullptr") { #if HT_AMD - HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-45"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130"); return; #endif DriverContext ctx; SECTION("Null array") { HIP_ARRAY_DESCRIPTOR desc; - desc.Format = formats[0]; + desc.Format = driverFormats[0]; desc.NumChannels = 4; desc.Width = 1024; desc.Height = 1024; @@ -374,7 +344,7 @@ TEST_CASE("Unit_hipArrayCreate_Nullptr") { TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") { DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; - desc.Format = GENERATE(from_range(std::begin(formats), std::end(formats))); + desc.Format = GENERATE(from_range(std::begin(driverFormats), std::end(driverFormats))); desc.NumChannels = GENERATE(-1, 0, 3, 5, 8); desc.Width = 1024; desc.Height = GENERATE(0, 1024); @@ -393,9 +363,9 @@ TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") { // create a bad format desc.Format = - std::accumulate(std::begin(formats), std::end(formats), formats[0], + std::accumulate(std::begin(driverFormats), std::end(driverFormats), driverFormats[0], [](auto i, auto f) { return static_cast(i + f); }); - for (auto&& format : formats) { + for (auto&& format : driverFormats) { REQUIRE(desc.Format != format); } diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index be41740242..1a8ec2f0d0 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -33,6 +33,7 @@ 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. @@ -67,53 +68,6 @@ static void Malloc3DArray_DiffSizes(int gpu) { REQUIRE_THREAD(pavail == avail); } -/* - * 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(); - 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); - } -} - TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") { Malloc3DArray_DiffSizes(0); HIP_CHECK_THREAD_FINALIZE(); @@ -173,7 +127,7 @@ void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected REQUIRE(expected_flags == queried_flags); #endif } -} +} // namespace TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, float2, float4) { @@ -187,15 +141,9 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, constexpr size_t size = 64; hipExtent extent; - SECTION("1D Array") { - extent = make_hipExtent(size, 0, 0); - } - SECTION("2D Array") { - extent = make_hipExtent(size, size, 0); - } - SECTION("3D Array") { - extent = make_hipExtent(size, size, size); - } + SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); } + SECTION("2D Array") { extent = make_hipExtent(size, size, 0); } + SECTION("3D Array") { extent = make_hipExtent(size, size, size); } HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); checkArrayIsExpected(array, desc, extent, flags); @@ -272,3 +220,203 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, us HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flag), hipErrorInvalidValue); } } + + +#if HT_AMD +constexpr std::array validFlags{hipArrayDefault}; +#else +constexpr std::array validFlags{ + hipArrayDefault, + hipArrayDefault | hipArraySurfaceLoadStore, + hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayCubemap, + hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayTextureGather}; +#endif + +hipExtent makeExtent(unsigned int flag, size_t s) { + if (flag == hipArrayTextureGather) { + return make_hipExtent(s, s, 0); + } + return make_hipExtent(s, s, s); +} + + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipMalloc3DArray_Negative_NullArrayPtr") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + constexpr size_t s = 6; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(nullptr, &desc, makeExtent(flag, s), flag), + hipErrorInvalidValue); +} + +// Providing the description pointer as nullptr should return an error +TEST_CASE("Unit_hipMalloc3DArray_Negative_NullDescPtr") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, nullptr, makeExtent(flag, s), flag), + hipErrorInvalidValue); +} + +// Zero width arrays are not allowed +TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroWidth") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(0, s, s), flag), + hipErrorInvalidValue); +} + +// Zero height arrays are only allowed for 1D arrays and layered arrays +TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroHeight") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array exceptions{hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore}; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + if (std::find(std::begin(exceptions), std::end(exceptions), flag) == std::end(exceptions)) { + // flag is not in list of exceptions + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(s, 0, s), flag), + hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + +#if HT_AMD + const unsigned int flag = 0xDEADBEEF; +#else + const unsigned int flag = + GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore, + hipArrayTextureGather | hipArrayCubemap, + hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap); +#endif + + CAPTURE(flag); + + REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), flag) == std::end(validFlags)); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), hipErrorInvalidValue); +} + +void testInvalidDescription(hipChannelFormatDesc desc){ + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + +#if HT_NVIDIA + hipError_t expectedError = hipErrorUnknown; +#else + hipError_t expectedError = hipErrorInvalidValue; +#endif + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), expectedError); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFormat") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + desc.f = GENERATE(hipChannelFormatKindNone, 0xBEEF); + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelLayout") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); + 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(bits, bits, bits, 0, formatKind), + 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); + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_8BitFloat") { + hipChannelFormatDesc desc = GENERATE(hipCreateChannelDesc(8, 0, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 8, 8, hipChannelFormatKindFloat)); + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_DifferentChannelSizes") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); + 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); + + INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX + << ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW); + + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelSize") { + const int badBits = GENERATE(-1, 0, 10, 100); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + hipChannelFormatDesc desc = hipCreateChannelDesc(badBits, badBits, badBits, badBits, formatKind); + + INFO("Number of bits: " << badBits); + + testInvalidDescription(desc); +} + + +// hipMalloc3DArray should handle the max numeric value gracefully. +TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") { + hipArray_t arrayPtr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + size_t size = std::numeric_limits::max(); + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag), + hipErrorInvalidValue); +} diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index 820e917313..f1f6ebc9d2 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -135,20 +135,6 @@ template size_t getAllocSize(const size_t width, const size_t heigh return sizeof(T) * width * (height ? height : 1); } - -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. @@ -409,7 +395,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho // pointer to the array in device memory hipArray_t arrayPtr{}; size_t width = 1024; - size_t height; + size_t height{}; SECTION("hipArrayDefault") { height = GENERATE(0, 1024); @@ -509,7 +495,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush // Arrays with channels of different size are not allowed. TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") { #if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); return; #endif const int bitsX = GENERATE(8, 16, 32); @@ -564,10 +550,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_ZeroWidth") { // 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); @@ -575,10 +557,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") { // 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); @@ -586,10 +564,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") { // 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; @@ -687,7 +661,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_3ChannelElement") { // 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"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); return; #endif const int bits = GENERATE(8, 16, 32);