EXSWCPHIPT-95 - More comprehensive tests for hipArrayCreate (#2702)

[ROCm/hip-tests commit: 5628a7c009]

projects/hip-tests/catch/unit/memory/CMakeLists.txt

@@ -1,4 +1,4 @@
-# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved.
+# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -21,6 +21,7 @@
 # Common Tests - Test independent of all platforms
 if(HIP_PLATFORM MATCHES "amd")
 set(TEST_SRC
+    DriverContext.cc
     memset.cc
     malloc.cc
     hipMemcpy2DToArray.cc
@@ -88,6 +89,7 @@ set(TEST_SRC
 )
 else()
 set(TEST_SRC
+    DriverContext.cc
     memset.cc
     malloc.cc
     hipMemcpy2DToArray.cc
@@ -159,4 +161,5 @@ endif()
 
 hip_add_exe_to_target(NAME MemoryTest
                       TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests)
+                      TEST_TARGET_NAME build_tests
+                      COMPILE_OPTIONS -std=c++14)

projects/hip-tests/catch/unit/memory/DriverContext.cc

@@ -0,0 +1,40 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "DriverContext.hh"
+#include <hip_test_common.hh>
+
+DriverContext::DriverContext() {
+#if HT_NVIDIA
+  HIP_CHECK(hipInit(0));
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  HIP_CHECK(hipDevicePrimaryCtxRetain(&ctx, device));
+  HIP_CHECK(hipCtxPushCurrent(ctx));
+#endif
+}
+
+DriverContext::~DriverContext() {
+#if HT_NVIDIA
+  HIP_CHECK(hipCtxPopCurrent(&ctx));
+  HIP_CHECK(hipDevicePrimaryCtxRelease(device));
+#endif
+}

projects/hip-tests/catch/unit/memory/DriverContext.hh

@@ -0,0 +1,41 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <hip_test_context.hh>
+
+class DriverContext {
+ private:
+#if HT_NVIDIA
+  hipCtx_t ctx;
+  hipDevice_t device;
+#endif
+
+ public:
+  DriverContext();
+  ~DriverContext();
+
+  // Rule of three
+  DriverContext(const DriverContext& other) = delete;
+  DriverContext(DriverContext&& other) noexcept = delete;
+};

projects/hip-tests/catch/unit/memory/hipArrayCommon.hh

@@ -0,0 +1,124 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+
+constexpr size_t BlockSize = 16;
+
+template <class T, size_t N, hipArray_Format Format> struct type_and_size_and_format {
+  using type = T;
+  static constexpr size_t size = N;
+  static constexpr hipArray_Format format = Format;
+};
+
+// Create a map of type to scalar type, vector size and scalar type format enum.
+// This is useful for creating simpler function that depend on the vector size.
+template <typename T> struct vector_info;
+template <>
+struct vector_info<int> : type_and_size_and_format<int, 1, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float> : type_and_size_and_format<float, 1, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short> : type_and_size_and_format<short, 1, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char> : type_and_size_and_format<char, 1, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<unsigned int>
+    : type_and_size_and_format<unsigned int, 1, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<unsigned short>
+    : type_and_size_and_format<unsigned short, 1, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<unsigned char>
+    : type_and_size_and_format<unsigned char, 1, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int2> : type_and_size_and_format<int, 2, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float2> : type_and_size_and_format<float, 2, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short2> : type_and_size_and_format<short, 2, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char2> : type_and_size_and_format<char, 2, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint2>
+    : type_and_size_and_format<unsigned int, 2, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort2>
+    : type_and_size_and_format<unsigned short, 2, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar2>
+    : type_and_size_and_format<unsigned char, 2, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+template <>
+struct vector_info<int4> : type_and_size_and_format<int, 4, HIP_AD_FORMAT_SIGNED_INT32> {};
+template <> struct vector_info<float4> : type_and_size_and_format<float, 4, HIP_AD_FORMAT_FLOAT> {};
+template <>
+struct vector_info<short4> : type_and_size_and_format<short, 4, HIP_AD_FORMAT_SIGNED_INT16> {};
+template <>
+struct vector_info<char4> : type_and_size_and_format<char, 4, HIP_AD_FORMAT_SIGNED_INT8> {};
+template <>
+struct vector_info<uint4>
+    : type_and_size_and_format<unsigned int, 4, HIP_AD_FORMAT_UNSIGNED_INT32> {};
+template <>
+struct vector_info<ushort4>
+    : type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
+template <>
+struct vector_info<uchar4>
+    : type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
+
+// read from a texture using normalized coordinates
+constexpr size_t ChannelToRead = 1;
+template <typename T>
+__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
+                                bool textureGather) {
+  // Calculate normalized texture coordinates
+  const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
+  const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
+  const float u = x / (float)width;
+
+  // Read from texture and write to global memory
+  if (height == 0) {
+    output[x] = tex1D<T>(texObj, u);
+  } else {
+    const float v = y / (float)height;
+    output[y * width + x] =
+        textureGather ? tex2Dgather<T>(texObj, u, v, ChannelToRead) : tex2D<T>(texObj, u, v);
+  }
+}
+
+template <typename T> void checkDataIsAscending(const std::vector<T>& hostData) {
+  bool allMatch = true;
+  size_t i = 0;
+  for (; i < hostData.size(); ++i) {
+    allMatch = allMatch && hostData[i] == static_cast<T>(i);
+    if (!allMatch) break;
+  }
+  INFO("hostData[" << i << "] == " << static_cast<T>(hostData[i]));
+  REQUIRE(allMatch);
+}
+
+inline size_t getFreeMem() {
+  size_t free = 0, total = 0;
+  HIP_CHECK(hipMemGetInfo(&free, &total));
+  return free;
+}

projects/hip-tests/catch/unit/memory/hipArrayCreate.cc

@@ -24,7 +24,11 @@ hipArrayCreate API test scenarios
 3. Multithreaded scenario
 */
 
+#include <array>
+#include <numeric>
 #include <hip_test_common.hh>
+#include "hipArrayCommon.hh"
+#include "DriverContext.hh"
 
 static constexpr auto NUM_W{4};
 static constexpr auto BIGNUM_W{100};
@@ -48,76 +52,31 @@ static constexpr auto ARRAY_LOOP{100};
 
 static void ArrayCreate_DiffSizes(int gpu) {
   HIP_CHECK(hipSetDevice(gpu));
-  std::vector<size_t> array_size;
-  array_size.push_back(NUM_W);
-  array_size.push_back(BIGNUM_W);
-  for (auto &size : array_size) {
-    HIP_ARRAY array[ARRAY_LOOP];
-    size_t tot, avail, ptot, pavail;
-    HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  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_ARRAY_DESCRIPTOR desc;
-      desc.NumChannels = 1;
-      if (size == NUM_W) {
-        desc.Width = NUM_W;
-        desc.Height = NUM_H;
-      } else {
-        desc.Width = BIGNUM_W;
-        desc.Height = BIGNUM_H;
-      }
-      desc.Format = HIP_AD_FORMAT_FLOAT;
       HIP_CHECK(hipArrayCreate(&array[i], &desc));
     }
     for (int i = 0; i < ARRAY_LOOP; i++) {
-      ARRAY_DESTROY(array[i]);
+      HIP_CHECK(hipArrayDestroy(array[i]));
     }
-    HIP_CHECK(hipMemGetInfo(&avail, &tot));
-    if ((pavail != avail)) {
+    const size_t avail = getFreeMem();
+    if (pavail != avail) {
       HIPASSERT(false);
     }
   }
 }
 
-/*Thread function*/
-static void ArrayCreateThreadFunc(int gpu) {
-  ArrayCreate_DiffSizes(gpu);
-}
-
 /* 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); }
 
-
-/* This testcase verifies the negative scenarios of
- * hipArrayCreate API
- */
-TEST_CASE("Unit_hipArrayCreate_Negative") {
-  HIP_ARRAY_DESCRIPTOR desc;
-  HIP_ARRAY array;
-  desc.Format = HIP_AD_FORMAT_FLOAT;
-  desc.NumChannels = 1;
-  desc.Width = NUM_W;
-  desc.Height = NUM_H;
-#if HT_NVIDIA
-  SECTION("NullPointer to Array") {
-    REQUIRE(hipArrayCreate(nullptr, &desc) != hipSuccess);
-  }
-
-  SECTION("NullPointer to Channel Descriptor") {
-    REQUIRE(hipArrayCreate(&array, nullptr) != hipSuccess);
-  }
-#endif
-  SECTION("Width 0 for Array Descriptor") {
-    desc.Width = 0;
-    REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess);
-  }
-
-  SECTION("Invalid NumChannels") {
-    desc.NumChannels = 3;
-    REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess);
-  }
-}
 /*
 This testcase verifies the hipArrayCreate API in multithreaded
 scenario by launching threads in parallel on multiple GPUs
@@ -129,16 +88,16 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") {
 
   devCnt = HipTest::getDeviceCount();
 
-  size_t tot, avail, ptot, pavail;
-  HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
+  const size_t pavail = getFreeMem();
   for (int i = 0; i < devCnt; i++) {
-    threadlist.push_back(std::thread(ArrayCreateThreadFunc, 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) {
+  for (auto& t : threadlist) {
     t.join();
   }
-  HIP_CHECK(hipMemGetInfo(&avail, &tot));
+  const size_t avail = getFreeMem();
 
   if (pavail != avail) {
     WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
@@ -146,3 +105,305 @@ 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
+constexpr auto MemoryTypeHost = hipMemoryTypeHost;
+constexpr auto MemoryTypeArray = hipMemoryTypeArray;
+constexpr auto NORMALIZED_COORDINATES = HIP_TRSF_NORMALIZED_COORDINATES;
+constexpr auto READ_AS_INTEGER = HIP_TRSF_READ_AS_INTEGER;
+#else
+constexpr auto MemoryTypeHost = CU_MEMORYTYPE_HOST;
+constexpr auto MemoryTypeArray = CU_MEMORYTYPE_ARRAY;
+// (EXSWCPHIPT-92) HIP equivalents not defined for CUDA backend.
+constexpr auto NORMALIZED_COORDINATES = CU_TRSF_NORMALIZED_COORDINATES;
+constexpr auto READ_AS_INTEGER = CU_TRSF_READ_AS_INTEGER;
+#endif
+
+// Copy data from host to the hiparray, accounting 1D or 2D arrays
+template <typename T>
+void copyToArray(hiparray dst, const std::vector<T>& src, const size_t height) {
+  const auto sizeInBytes = src.size() * sizeof(T);
+  if (height == 0) {
+    // FIXME(EXSWCPHIPT-64) remove cast when API is fixed (will require major version change)
+    HIP_CHECK(hipMemcpyHtoA(reinterpret_cast<hipArray*>(dst), 0, src.data(), sizeInBytes));
+  } else {
+    const auto pitch = sizeInBytes / height;
+    hip_Memcpy2D copyParams{};
+    copyParams.srcMemoryType = MemoryTypeHost;
+    copyParams.srcXInBytes = 0;  // x offset
+    copyParams.srcY = 0;         // y offset
+    copyParams.srcHost = src.data();
+    copyParams.srcPitch = pitch;
+
+
+    copyParams.dstMemoryType = MemoryTypeArray;
+    copyParams.dstXInBytes = 0;  // x offset
+    copyParams.dstY = 0;         // y offset
+    copyParams.dstArray = dst;
+
+    copyParams.WidthInBytes = pitch;
+    copyParams.Height = height;
+
+    HIP_CHECK(hipMemcpyParam2D(&copyParams));
+  }
+}
+
+// Test the allocated array by generating a texture from it then reading from that texture.
+// Textures are read-only, so write to the array then copy that into normal device memory.
+template <typename T>
+void testArrayAsTexture(hiparray array, const size_t width, const size_t height) {
+  using vec_info = vector_info<T>;
+  using scalar_type = typename vec_info::type;
+  const auto h = height ? height : 1;
+  const auto size = sizeof(T) * width * h;
+
+  // set hip array
+  std::vector<scalar_type> hostData(width * h * vec_info::size);
+  // assigned ascending values to the data array to show indexing is working
+  std::iota(std::begin(hostData), std::end(hostData), 0);
+
+  copyToArray(array, hostData, height);
+
+  // create texture
+  hipTextureObject_t textObj{};
+
+  HIP_RESOURCE_DESC resDesc{};
+  memset(&resDesc, 0, sizeof(HIP_RESOURCE_DESC));
+  resDesc.resType = HIP_RESOURCE_TYPE_ARRAY;
+  resDesc.res.array.hArray = array;
+  resDesc.flags = 0;
+
+  HIP_TEXTURE_DESC texDesc{};
+  memset(&texDesc, 0, sizeof(HIP_TEXTURE_DESC));
+  // use the actual values in the texture, not normalized data
+  texDesc.filterMode = HIP_TR_FILTER_MODE_POINT;
+  // Use normalized coordinates and also read the data in the original data type
+  texDesc.flags |= NORMALIZED_COORDINATES | READ_AS_INTEGER;
+
+  HIP_CHECK(hipTexObjectCreate(&textObj, &resDesc, &texDesc, nullptr));
+
+  // run kernel
+  T* device_data{};
+  HIP_CHECK(hipMalloc(&device_data, size));
+  readFromTexture<<<dim3(width / BlockSize, height ? height / BlockSize : 1, 1),
+                    dim3(BlockSize, height ? BlockSize : 1, 1)>>>(device_data, textObj, width,
+                                                                  height, false);
+  HIP_CHECK(hipGetLastError());  // check for errors when running the kernel
+
+  // copy data back and then test it
+  std::fill(std::begin(hostData), std::end(hostData), 0);
+  HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost));
+
+  checkDataIsAscending(hostData);
+
+  // clean up
+  HIP_CHECK(hipTexObjectDestroy(textObj));
+  HIP_CHECK(hipFree(device_data));
+}
+
+// Selection of types chosen since trying all types would be slow to compile
+// Test the happy path of the hipArrayCreate
+TEMPLATE_TEST_CASE("Unit_hipArrayCreate_happy", "", uint, int, int4, ushort, short2, char, uchar2,
+                   char4, float, float2, float4) {
+#if HT_AMD
+  if (std::is_same<uint, TestType>::value || std::is_same<ushort, TestType>::value ||
+      std::is_same<uchar2, TestType>::value) {
+    HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-62");
+    return;
+  }
+#endif
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+  desc.Width = 1024;
+  desc.Height = GENERATE(0, 1024);
+
+  size_t initFree = getFreeMem();
+
+  // pointer to the array in device memory
+  hiparray array{};
+
+  HIP_CHECK(hipArrayCreate(&array, &desc));
+
+  testArrayAsTexture<TestType>(array, desc.Width, desc.Height);
+
+  size_t finalFree = getFreeMem();
+
+  const size_t allocSize = sizeof(TestType) * desc.Width * (desc.Height ? desc.Height : 1);
+  // will be aligned to some size, so this is not exact
+  REQUIRE(initFree - finalFree >= allocSize);
+
+  HIP_CHECK(hipArrayDestroy(array));
+}
+
+
+// Only widths and Heights up to the maxTexture size is supported
+TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort, short2, char,
+                   uchar2, char4, float, float2, float4) {
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+
+  int device;
+  HIP_CHECK(hipGetDevice(&device));
+  hipDeviceProp_t prop;
+  HIP_CHECK(hipGetDeviceProperties(&prop, device));
+
+  hiparray array{};
+  SECTION("Happy") {
+    SECTION("1D - Max") {
+      desc.Width = prop.maxTexture1D;
+      desc.Height = 0;
+    }
+    SECTION("2D - Max Width") {
+      desc.Width = prop.maxTexture2D[0];
+      desc.Height = 64;
+    }
+    SECTION("2D - Max Height") {
+      desc.Width = 64;
+      desc.Height = prop.maxTexture2D[1];
+    }
+    SECTION("2D - Max Width and Height") {
+      desc.Width = prop.maxTexture2D[0];
+      desc.Height = prop.maxTexture2D[1];
+    }
+    auto maxArrayCreateError = hipArrayCreate(&array, &desc);
+    // this can try to alloc many GB of memory, so out of memory is acceptable
+    // return to avoid destroy
+    if (maxArrayCreateError == hipErrorOutOfMemory) return;
+    HIP_CHECK(maxArrayCreateError);
+    HIP_CHECK(hipArrayDestroy(array));
+  }
+  SECTION("Negative") {
+    SECTION("1D - More Than Max") {
+      desc.Width = prop.maxTexture1D + 1;
+      desc.Height = 0;
+    }
+    SECTION("2D - More Than Max Width") {
+      desc.Width = prop.maxTexture2D[0] + 1;
+      desc.Height = 64;
+    }
+    SECTION("2D - More Than Max Height") {
+      desc.Width = 64;
+      desc.Height = prop.maxTexture2D[1] + 1;
+    }
+    SECTION("2D - More Than Max Width and Height") {
+      desc.Width = prop.maxTexture2D[0] + 1;
+      desc.Height = prop.maxTexture2D[1] + 1;
+    }
+    HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+  }
+}
+
+// zero-width array is not supported
+TEST_CASE("Unit_hipArrayCreate_ZeroWidth") {
+  DriverContext ctx;
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = formats[0];
+  desc.NumChannels = 4;
+  desc.Width = 0;
+  desc.Height = GENERATE(0, 1024);
+
+  // pointer to the array in device memory
+  hiparray array;
+  HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+}
+
+// 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");
+  return;
+#endif
+  DriverContext ctx;
+  SECTION("Null array") {
+    HIP_ARRAY_DESCRIPTOR desc;
+    desc.Format = formats[0];
+    desc.NumChannels = 4;
+    desc.Width = 1024;
+    desc.Height = 1024;
+
+    HIP_CHECK_ERROR(hipArrayCreate(nullptr, &desc), hipErrorInvalidValue);
+  }
+  SECTION("Null Description") {
+    hiparray array;
+    HIP_CHECK_ERROR(hipArrayCreate(&array, nullptr), hipErrorInvalidValue);
+  }
+}
+
+// Only elements with 1,2, or 4 channels is supported
+TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") {
+  DriverContext ctx;
+  HIP_ARRAY_DESCRIPTOR desc;
+  desc.Format = GENERATE(from_range(std::begin(formats), std::end(formats)));
+  desc.NumChannels = GENERATE(-1, 0, 3, 5, 8);
+  desc.Width = 1024;
+  desc.Height = GENERATE(0, 1024);
+
+  hiparray array;
+
+  INFO("Format: " << formatToString(desc.Format) << " NumChannels: " << desc.NumChannels
+                  << " Height: " << desc.Height)
+  HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+}
+
+// Only certain channel formats are acceptable.
+TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") {
+  DriverContext ctx;
+  HIP_ARRAY_DESCRIPTOR desc;
+
+  // create a bad format
+  desc.Format =
+      std::accumulate(std::begin(formats), std::end(formats), formats[0],
+                      [](auto i, auto f) { return static_cast<decltype(desc.Format)>(i + f); });
+  for (auto&& format : formats) {
+    REQUIRE(desc.Format != format);
+  }
+
+  desc.NumChannels = 4;
+  desc.Width = 1024;
+  desc.Height = GENERATE(0, 1024);
+
+  hiparray array;
+
+  INFO("Format: " << formatToString(desc.Format) << " Height: " << desc.Height)
+  HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue);
+}

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

@@ -27,9 +27,8 @@ hipMallocArray API test scenarios
 
 #include <hip_test_common.hh>
 #include <limits>
-#if defined(_WIN32) || defined(_WIN64)
 #include <numeric>
-#endif
+#include "hipArrayCommon.hh"
 
 static constexpr auto NUM_W{4};
 static constexpr auto BIGNUM_W{100};
@@ -86,7 +85,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.
+    // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken.
     threadlist.push_back(std::thread(MallocArray_DiffSizes, i));
   }
 
@@ -101,63 +100,8 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
   }
 }
 
-
-constexpr size_t BlockSize = 16;
-
-template <class T, size_t N> struct type_and_size {
-  using type = T;
-  static constexpr size_t size = N;
-};
-
-// scalars are interpreted as a vector of 1 length.
-// template <size_t N> using int_constant = std::integral_constant<size_t, N>;
-template <typename T> struct vector_info;
-template <> struct vector_info<int> : type_and_size<int, 1> {};
-template <> struct vector_info<float> : type_and_size<float, 1> {};
-template <> struct vector_info<short> : type_and_size<short, 1> {};
-template <> struct vector_info<char> : type_and_size<char, 1> {};
-template <> struct vector_info<unsigned int> : type_and_size<unsigned int, 1> {};
-template <> struct vector_info<unsigned short> : type_and_size<unsigned short, 1> {};
-template <> struct vector_info<unsigned char> : type_and_size<unsigned char, 1> {};
-
-template <> struct vector_info<int2> : type_and_size<int, 2> {};
-template <> struct vector_info<float2> : type_and_size<float, 2> {};
-template <> struct vector_info<short2> : type_and_size<short, 2> {};
-template <> struct vector_info<char2> : type_and_size<char, 2> {};
-template <> struct vector_info<uint2> : type_and_size<unsigned int, 2> {};
-template <> struct vector_info<ushort2> : type_and_size<unsigned short, 2> {};
-template <> struct vector_info<uchar2> : type_and_size<unsigned char, 2> {};
-
-template <> struct vector_info<int4> : type_and_size<int, 4> {};
-template <> struct vector_info<float4> : type_and_size<float, 4> {};
-template <> struct vector_info<short4> : type_and_size<short, 4> {};
-template <> struct vector_info<char4> : type_and_size<char, 4> {};
-template <> struct vector_info<uint4> : type_and_size<unsigned int, 4> {};
-template <> struct vector_info<ushort4> : type_and_size<unsigned short, 4> {};
-template <> struct vector_info<uchar4> : type_and_size<unsigned char, 4> {};
-
 // Kernels ///////////////////////////////////////
 
-// read from a texture using normalized coordinates
-constexpr size_t ChannelToRead = 1;
-template <typename T>
-__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
-                                bool textureGather) {
-  // Calculate normalized texture coordinates
-  const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
-  const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
-  const float u = x / (float)width;
-
-  // Read from texture and write to global memory
-  if (height == 0) {
-    output[x] = tex1D<T>(texObj, u);
-  } else {
-    const float v = y / (float)height;
-    output[y * width + x] =
-        textureGather ? tex2Dgather<T>(texObj, u, v, ChannelToRead) : tex2D<T>(texObj, u, v);
-  }
-}
-
 template <typename T> __device__ void addOne(T* a) {
   using scalar_type = typename vector_info<T>::type;
   auto as = reinterpret_cast<scalar_type*>(a);
@@ -190,16 +134,6 @@ template <typename T> size_t getAllocSize(const size_t width, const size_t heigh
   return sizeof(T) * width * (height ? height : 1);
 }
 
-template <typename T> void checkDataIsAscending(const std::vector<T>& hostData) {
-  bool allMatch = true;
-  size_t i = 0;
-  for (; i < hostData.size(); ++i) {
-    allMatch = allMatch && hostData[i] == static_cast<T>(i);
-    if (!allMatch) break;
-  }
-  INFO("hostData[" << i << "] == " << static_cast<T>(hostData[i]));
-  REQUIRE(allMatch);
-}
 
 const char* channelFormatString(hipChannelFormatKind formatKind) noexcept {
   switch (formatKind) {
@@ -458,12 +392,6 @@ void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t he
   HIP_CHECK(hipFree(device_data));
 }
 
-size_t getFreeMem() {
-  size_t free = 0, total = 0;
-  HIP_CHECK(hipMemGetInfo(&free, &total));
-  return free;
-}
-
 // The happy path of a default array and a SurfaceLoadStore array should work
 // Selection of types chosen to reduce compile times
 TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, short2, char, uchar2,
@@ -526,6 +454,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush
   HIP_CHECK(hipGetDevice(&device));
   hipDeviceProp_t prop;
   HIP_CHECK(hipGetDeviceProperties(&prop, device));
+
   size_t width, height;
   hipArray_t array{};
   hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
@@ -549,7 +478,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush
       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
+    // this can try to alloc many GB of memory, so out of memory is acceptable
     if (maxArrayCreateError == hipErrorOutOfMemory) return;
     HIP_CHECK(maxArrayCreateError);
     HIP_CHECK(hipFreeArray(array));