EXSWCPHIPT-134 - Negative tests for hipMalloc3DArray and hipArray3DCreate (#2784)

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 <limits>
 #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<TestType>;
   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<unsigned int, 9> validFlags{
+    0,
+    hipArraySurfaceLoadStore,
+    hipArrayLayered,
+    hipArrayLayered | hipArraySurfaceLoadStore,
+    hipArrayCubemap,
+    hipArrayCubemap | hipArrayLayered,
+    hipArrayCubemap | hipArraySurfaceLoadStore,
+    hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore,
+    hipArrayTextureGather};
+#else
+constexpr std::array<unsigned int, 5> 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<unsigned int, 2> exceptions{CUDA_ARRAY3D_LAYERED,
+                                         CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST};
+#else
+  std::array<unsigned int, 0> 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<hipArray_Format>(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<size_t>::max());
+
+  testInvalidDescription(desc);
+}

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<hipArray_Format> 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";
+  }
+}

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<hipArray_Format> 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<decltype(desc.Format)>(i + f); });
-  for (auto&& format : formats) {
+  for (auto&& format : driverFormats) {
     REQUIRE(desc.Format != format);
   }
 

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<float>();
-  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<int>::max(), std::numeric_limits<int>::max(),
-                               std::numeric_limits<int>::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<unsigned int, 1> validFlags{hipArrayDefault};
+#else
+constexpr std::array<unsigned int, 9> 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<float4>();
+  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<float4>();
+
+  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<float4>();
+  std::array<unsigned int, 2> 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<float4>();
+
+#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<float4>();
+  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<float>();
+
+  size_t size = std::numeric_limits<size_t>::max();
+  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag),
+                  hipErrorInvalidValue);
+}

tests/catch/unit/memory/hipMallocArray.cc

@@ -135,20 +135,6 @@ template <typename T> 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<float4>();
 
   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<float4>();
 
   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);