EXSWCPHIPT-113 - Unit tests for hipArray3DCreate for default and surface arrays (#2714)

[ROCm/hip-tests commit: 646eaa43ea]
2022-07-11 11:41:32 +01:00
@@ -82,6 +82,7 @@ set(TEST_SRC
    hipMalloc3D.cc
    hipMalloc3DArray.cc
    hipArrayCreate.cc
+    hipArray3DCreate.cc
    hipDrvMemcpy3D.cc
    hipDrvMemcpy3DAsync.cc
    hipPointerGetAttribute.cc
@@ -150,6 +151,7 @@ set(TEST_SRC
    hipMalloc3D.cc
    hipMalloc3DArray.cc
    hipArrayCreate.cc
+    hipArray3DCreate.cc
    hipDrvMemcpy3D.cc
    hipDrvMemcpy3DAsync.cc
    hipPointerGetAttribute.cc
@@ -24,17 +24,13 @@ THE SOFTWARE.
 #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
 }
@@ -26,10 +26,8 @@ THE SOFTWARE.

 class DriverContext {
 private:
-#if HT_NVIDIA
  hipCtx_t ctx;
  hipDevice_t device;
-#endif

 public:
  DriverContext();
@@ -0,0 +1,166 @@
+/*
+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 "hipArrayCommon.hh"
+
+namespace {
+void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) {
+// hipArray3DGetDescriptor doesn't currently exist (EXSWCPHIPT-87)
+#if HT_AMD
+  std::ignore = array;
+  std::ignore = expected_desc;
+#else
+  CUDA_ARRAY3D_DESCRIPTOR queried_desc;
+  cuArray3DGetDescriptor(&queried_desc, array);
+
+  REQUIRE(queried_desc.Width == expected_desc.Width);
+  REQUIRE(queried_desc.Height == expected_desc.Height);
+  REQUIRE(queried_desc.Depth == expected_desc.Depth);
+  REQUIRE(queried_desc.Format == expected_desc.Format);
+  REQUIRE(queried_desc.NumChannels == expected_desc.NumChannels);
+  REQUIRE(queried_desc.Flags == expected_desc.Flags);
+#endif
+}
+}  // namespace
+
+TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float,
+                   float2, float4) {
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  hiparray array;
+  HIP_ARRAY3D_DESCRIPTOR desc{};
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+#if HT_AMD
+  desc.Flags = 0;
+#else
+  desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST);
+#endif
+
+  constexpr size_t size = 64;
+
+  std::vector<hipExtent> extents{
+      {size, 0, 0},       // 1D array
+      {size, size, 0},    // 2D array
+      {size, size, size}  // 3D array
+  };
+
+  for (auto& extent : extents) {
+    desc.Width = extent.width;
+    desc.Height = extent.height;
+    desc.Depth = extent.depth;
+
+    CAPTURE(desc.Width, desc.Height, desc.Depth);
+
+    HIP_CHECK(hipArray3DCreate(&array, &desc));
+    checkArrayIsExpected(array, desc);
+    HIP_CHECK(hipArrayDestroy(array));
+  }
+}
+
+TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, ushort2,
+                   unsigned char, float, float4) {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97");
+  return;
+#endif
+
+  using vec_info = vector_info<TestType>;
+  DriverContext ctx;
+
+  hiparray array;
+  HIP_ARRAY3D_DESCRIPTOR desc{};
+  desc.Format = vec_info::format;
+  desc.NumChannels = vec_info::size;
+#if HT_AMD
+  desc.Flags = 0;
+#else
+  desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST);
+  if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) {
+    HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58");
+    return;
+  }
+#endif
+  CAPTURE(desc.Flags);
+
+  const Sizes sizes(desc.Flags);
+  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)));
+
+    desc.Width = extent.width;
+    desc.Height = extent.height;
+    desc.Depth = extent.depth;
+
+    CAPTURE(desc.Width, desc.Height, desc.Depth);
+
+    auto maxArrayCreateError = hipArray3DCreate(&array, &desc);
+    // 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);
+    HIP_CHECK(hipArrayDestroy(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)));
+
+    desc.Width = extent.width;
+    desc.Height = extent.height;
+    desc.Depth = extent.depth;
+
+    CAPTURE(desc.Width, desc.Height, desc.Depth);
+
+    HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue);
+  }
+}
@@ -131,8 +131,16 @@ struct Sizes {
  Sizes(unsigned int flag) {
    int device;
    HIP_CHECK(hipGetDevice(&device));
+    static_assert(
+        hipArrayDefault == 0,
+        "hipArrayDefault is assumed to be equivalent to 0 for the following switch statment");
+#if HT_NVIDIA
+    static_assert(hipArraySurfaceLoadStore == CUDA_ARRAY3D_SURFACE_LDST,
+                  "hipArraySurface is assumed to be equivalent to CUDA_ARRAY3D_SURFACE_LDST for "
+                  "the following switch statment");
+#endif
    switch (flag) {
-      case hipArrayDefault: {
+      case hipArrayDefault: {  // 0
        hipDeviceProp_t prop;
        HIP_CHECK(hipGetDeviceProperties(&prop, device));
        max1D = prop.maxTexture1D;
@@ -140,7 +148,7 @@ struct Sizes {
        max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]};
        return;
      }
-      case hipArraySurfaceLoadStore: {
+      case hipArraySurfaceLoadStore: {  // CUDA_ARRAY3D_SURFACE_LDST
        int value;
        HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device));
        max1D = value;
@@ -145,6 +145,7 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") {
  }
 }

+namespace {
 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)
@@ -172,6 +173,7 @@ void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected
  REQUIRE(expected_flags == queried_flags);
 #endif
 }
+}

 TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float,
                   float2, float4) {
@@ -187,19 +189,16 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4,

  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));
  }

+  HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags));
  checkArrayIsExpected(array, desc, extent, flags);
-
  HIP_CHECK(hipFreeArray(array));
 }