EXSWCPHIPT-112 - Unit test for hipMalloc3DArray for default and surface arrays (#2713)

[ROCm/hip commit: c0deb17bbc]

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<int, 2> max2D;
+  std::array<int, 3> 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;
+      }
+    }
+  }
+};

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<std::pair<size_t, size_t>> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}};
-  for (auto& size : array_size) {
-    std::array<HIP_ARRAY, ARRAY_LOOP> 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_t, size_t> size =
+      GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H));
+  std::array<HIP_ARRAY, ARRAY_LOOP> 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);
   }

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

@@ -25,9 +25,9 @@ hipMalloc3DArray API test scenarios
 4. Multithreaded scenario
 */
 
-
-
+#include <array>
 #include <hip_test_common.hh>
+#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<int> 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<float>();
+  std::array<hipArray_t, ARRAY_LOOP> 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<float>();
+  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<int>::max(),
-            std::numeric_limits<int>::max(),
-            std::numeric_limits<int>::max()),
-            hipArrayDefault) != hipSuccess);
+    REQUIRE(hipMalloc3DArray(
+                &arr, &channelDesc,
+                make_hipExtent(std::numeric_limits<int>::max(), std::numeric_limits<int>::max(),
+                               std::numeric_limits<int>::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<std::thread> 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<TestType>();
+#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<TestType>();
+#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<hipExtent> 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<hipExtent> 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);
+  }
+}

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

@@ -30,11 +30,11 @@ hipMallocArray API test scenarios
 #include <numeric>
 #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<std::pair<size_t, size_t>> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}};
-  for (auto& size : array_size) {
-    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++) {
-      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_t, size_t> size =
+      GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H));
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+  std::array<hipArray_t, ARRAY_LOOP> 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<std::thread> 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<TestType>();
   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);
   }