Added negative tests for hipMallocArray (#2699)

* Added negative tests for hipMallocArray

* fix numeric limit test for nvidia

[ROCm/hip commit: f6efba0b18]

projects/hip/tests/catch/unit/memory/hipMallocArray.cc

@@ -26,6 +26,7 @@ hipMallocArray API test scenarios
 */
 
 #include <hip_test_common.hh>
+#include <limits>
 
 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<size_t> array_size;
+  std::vector<std::pair<size_t, size_t>> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}};
-  array_size.push_back(NUM_W);
-  array_size.push_back(BIGNUM_W);
   for (auto& size : array_size) {
-    hipArray* A_d[ARRAY_LOOP];
+    std::array<hipArray_t, ARRAY_LOOP> A_d;
     size_t tot, avail, ptot, pavail;
     hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
     HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
     for (int i = 0; i < ARRAY_LOOP; i++) {
-      if (size == NUM_W) {
+      HIP_CHECK(
-        HIP_CHECK(hipMallocArray(&A_d[i], &desc, NUM_W, NUM_H, hipArrayDefault));
+          hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), 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]));
@@ -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<float>();
-#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<int>::max(),
-                           std::numeric_limits<int>::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 <typename T> void checkDataIsAscending(const std::vector<T>& 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<TestType>();
@@ -381,7 +357,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho
     HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault));
     testArrayAsTexture<TestType>(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<TestType>();
+  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<float4>();
+
+  // 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<float4>();
+
+  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<float4>();
+
+  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<hipChannelFormatKind>(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<float>();
+
+  size_t size = std::numeric_limits<size_t>::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);
+}