EXSWHTEC-100 - Miscellaneous modifications to existing memory tests

tests/catch/include/resource_guards.hh

@@ -80,17 +80,74 @@ template <typename T> class LinearAllocGuard {
     }
   }
 
-  T* ptr() { return ptr_; };
-  T* const ptr() const { return ptr_; };
-  T* host_ptr() { return host_ptr_; }
-  T* const host_ptr() const { return host_ptr(); }
+  T* ptr() const { return ptr_; };
+  T* host_ptr() const { return host_ptr_; }
 
  private:
   const LinearAllocs allocation_type_;
   T* ptr_ = nullptr;
   T* host_ptr_ = nullptr;
 };
 
+template <typename T> class LinearAllocGuardMultiDim {
+  protected:
+  LinearAllocGuardMultiDim(hipExtent extent)
+  : extent_{extent} {}
+
+  ~LinearAllocGuardMultiDim() {
+    static_cast<void>(hipFree(pitched_ptr_.ptr));
+  }
+
+  public:
+  T* ptr() const { return reinterpret_cast<T*>(pitched_ptr_.ptr); };
+
+  size_t pitch() const { return pitched_ptr_.pitch; }
+
+  hipExtent extent() const { return extent_; }
+
+  hipPitchedPtr pitched_ptr() const { return pitched_ptr_; }
+
+  size_t width() const { return extent_.width; }
+
+  size_t width_logical() const { return extent_.width / sizeof(T); }
+
+  size_t height() const { return extent_.height; }
+
+  public:
+  hipPitchedPtr pitched_ptr_;
+  const hipExtent extent_;
+};
+
+template <typename T> class LinearAllocGuard2D : public LinearAllocGuardMultiDim<T> {
+  public:
+  LinearAllocGuard2D(const size_t width_logical, const size_t height) 
+  : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, 1)}
+  {
+    HIP_CHECK(hipMallocPitch(&this->pitched_ptr_.ptr, &this->pitched_ptr_.pitch, this->extent_.width, this->extent_.height));
+  }
+
+  LinearAllocGuard2D(const LinearAllocGuard2D&) = delete;
+  LinearAllocGuard2D(LinearAllocGuard2D&&) = delete;
+};
+
+template <typename T> class LinearAllocGuard3D : public LinearAllocGuardMultiDim<T> {
+  public:
+  LinearAllocGuard3D(const size_t width_logical, const size_t height, const size_t depth)
+  : LinearAllocGuardMultiDim<T>{make_hipExtent(width_logical * sizeof(T), height, depth)}
+  {
+    HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
+  }
+
+  LinearAllocGuard3D(const hipExtent extent) : LinearAllocGuardMultiDim<T>(extent) {
+    HIP_CHECK(hipMalloc3D(&this->pitched_ptr_, this->extent_));
+  }
+
+  LinearAllocGuard3D(const LinearAllocGuard3D&) = delete;
+  LinearAllocGuard3D(LinearAllocGuard3D&&) = delete;
+
+  size_t depth() const { return this->extent_.depth; }
+};
+
 enum class Streams { nullstream, perThread, created };
 
 class StreamGuard {

tests/catch/include/utils.hh

@@ -54,6 +54,37 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele
   ArrayFindIfNot(array, array + num_elements, expected_value);
 }
 
+template <typename T, typename F>
+void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height,
+                         const size_t depth, F expected_value_generator) {
+  for (int z = 0; z < depth; ++z) {
+    for (int y = 0; y < height; ++y) {
+      for (int x = 0; x < width; ++x) {
+        const auto slice = reinterpret_cast<uint8_t*>(ptr) + pitch * height * z;
+        const auto row = slice + pitch * y;
+        if (reinterpret_cast<T*>(row)[x] != expected_value_generator(x, y, z)) {
+          INFO("Mismatch at indices: " << x << ", " << y << ", " << z);
+          REQUIRE(reinterpret_cast<T*>(row)[x] == expected_value_generator(x, y, z));
+        }
+      }
+    }
+  }
+}
+
+template <typename T, typename F>
+void PitchedMemorySet(T* const ptr, const size_t pitch, const size_t width, const size_t height,
+                      const size_t depth, F expected_value_generator) {
+  for (int z = 0; z < depth; ++z) {
+    for (int y = 0; y < height; ++y) {
+      for (int x = 0; x < width; ++x) {
+        const auto slice = reinterpret_cast<uint8_t*>(ptr) + pitch * height * z;
+        const auto row = slice + pitch * y;
+        reinterpret_cast<T*>(row)[x] = expected_value_generator(x, y, z);
+      }
+    }
+  }
+}
+
 template <typename T>
 __global__ void VectorIncrement(T* const vec, const T increment_value, size_t N) {
   size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
@@ -82,6 +113,18 @@ static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
   }
 }
 
+template <typename T>
+__global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) {
+  const auto x = blockIdx.x * blockDim.x + threadIdx.x;
+  const auto y = blockIdx.y * blockDim.y + threadIdx.y;
+  const auto z = blockIdx.z * blockDim.z + threadIdx.z;
+  if (x < w && y < h && z < d) {
+    char* const slice = reinterpret_cast<char*>(out) + pitch * h * z;
+    char* const row = slice + pitch * y;
+    reinterpret_cast<T*>(row)[x] = z * w * h + y * w + x;
+  }
+}
+
 inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
   int ticks_per_ms = 0;
   // Clock rate is in kHz => number of clock ticks in a millisecond

tests/catch/unit/memory/hipFree.cc

@@ -48,11 +48,10 @@ using namespace std::chrono_literals;
 const std::chrono::duration<uint64_t, std::milli> delay = 50ms;
 constexpr size_t numAllocs = 10;
 
-#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
-TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float4) {
-  TestType* devPtr{};
+TEST_CASE("Unit_hipFreeImplicitSyncDev") {
+  int* devPtr{};
   size_t size_mult = GENERATE(1, 32, 64, 128, 256);
-  HIP_CHECK(hipMalloc(&devPtr, sizeof(TestType) * size_mult));
+  HIP_CHECK(hipMalloc(&devPtr, sizeof(*devPtr) * size_mult));
 
   HipTest::runKernelForDuration(delay);
   // make sure device is busy
@@ -61,11 +60,11 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float
   HIP_CHECK(hipStreamQuery(nullptr));
 }
 
-TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, float4) {
-  TestType* hostPtr{};
+TEST_CASE("Unit_hipFreeImplicitSyncHost") {
+  int* hostPtr{};
   size_t size_mult = GENERATE(1, 32, 64, 128, 256);
 
-  HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(TestType) * size_mult));
+  HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(*hostPtr) * size_mult));
 
   HipTest::runKernelForDuration(delay);
   // make sure device is busy
@@ -74,7 +73,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, floa
   HIP_CHECK(hipStreamQuery(nullptr));
 }
 
-#if HT_NVIDIA  // Meaningless at the moment, since we are not running wait kernel on nvidia.
+#if HT_NVIDIA
 TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, float4) {
   using vec_info = vector_info<TestType>;
   DriverContext ctx;
@@ -134,7 +133,6 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo
   }
 }
 
-#endif
 #endif
 
 // Freeing a invalid pointer with on device
@@ -165,8 +163,6 @@ TEST_CASE("Unit_hipFreeNegativeHost") {
 #if HT_NVIDIA
 TEST_CASE("Unit_hipFreeNegativeArray") {
   DriverContext ctx;
-  hipArray_t arrayPtr{};
-  hiparray cuArrayPtr{};
 
   SECTION("ArrayFree") { HIP_CHECK(hipFreeArray(nullptr)); }
   SECTION("ArrayDestroy") {

tests/catch/unit/memory/hipHostGetDevicePointer.cc

@@ -21,11 +21,19 @@ THE SOFTWARE.
 */
 
 #include <hip_test_common.hh>
+#include <utils.hh>
 
 TEST_CASE("Unit_hipHostGetDevicePointer_Negative") {
   int* hPtr{nullptr};
+  int* dPtr{nullptr};
   HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int)));
 
+  if (!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) {
+    HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 0),
+                    hipErrorNotSupported);
+    return;
+  }
+
   SECTION("Nullptr as device") {
     HIP_CHECK_ERROR(hipHostGetDevicePointer(nullptr, hPtr, 0), hipErrorInvalidValue);
   }
@@ -36,13 +44,29 @@ TEST_CASE("Unit_hipHostGetDevicePointer_Negative") {
                     hipErrorInvalidValue);
   }
 
-  // Not adding check for flags since CUDA spec states that there might be more values added to it
+  SECTION("Non pinned memory as host") {
+    int* hPtr = reinterpret_cast<int*>(malloc(sizeof(*hPtr)));
+    HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 0),
+                    hipErrorInvalidValue);
+    free(hPtr);
+  }
+
+  SECTION("Flags non-zero") {
+    HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 1),
+                    hipErrorInvalidValue);
+  }
+
   HIP_CHECK(hipHostFree(hPtr));
 }
 
 template <typename T> __global__ void set(T* ptr, T val) { *ptr = val; }
 
 TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") {
+  if(!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) {
+    HipTest::HIP_SKIP_TEST("Device does not support mapping host memory"); 
+    return;
+  }
+
   int* hPtr{nullptr};
   HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int)));
 
@@ -71,8 +95,8 @@ TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") {
     HIP_CHECK(hipDeviceSynchronize());
     HIP_CHECK(hipHostUnregister(&res));
 
-    REQUIRE(value == 10);
+    REQUIRE(res == value);
   }
 
   HIP_CHECK(hipHostFree(hPtr));
-}
+}

tests/catch/unit/memory/hipHostRegister.cc

@@ -27,9 +27,10 @@ This testfile verifies the following scenarios of hipHostRegister API
 2. hipHostRegister and perform hipMemcpy on it.
 */
 
+#include "hip/hip_runtime_api.h"
 #include <hip_test_common.hh>
 #include <hip_test_helper.hh>
-#include "hip/hip_runtime_api.h"
+#include <utils.hh>
 
 #define OFFSET 128
 static constexpr auto LEN{1024 * 1024};
@@ -63,9 +64,7 @@ void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internal
   HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));
 
   // Make sure the copy worked
-  for (size_t i = 0; i < numElements; i++) {
-    REQUIRE(Bh[i] == A[i]);
-  }
+  ArrayMismatch(A, Bh, numElements);
 
   if (internalRegister) {
     HIP_CHECK(hipHostUnregister(A));

tests/catch/unit/memory/hipHostUnregister.cc

@@ -68,6 +68,12 @@ TEST_CASE("Unit_hipHostUnregister_NullPtr") {
   HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue);
 }
 
+TEST_CASE("Unit_hipHostUnregister_Ptr_Different_Than_Specified_To_Register") {
+  std::vector<int> alloc(2);
+  HIP_CHECK(hipHostRegister(alloc.data(), alloc.size(), 0));
+  HIP_CHECK_ERROR(hipHostUnregister(&alloc.data()[1]), hipErrorHostMemoryNotRegistered);
+}
+
 TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") {
   auto x = std::unique_ptr<int>(new int);
   HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered);

tests/catch/unit/memory/hipMallocPitch.cc

@@ -228,6 +228,21 @@ TEST_CASE("Unit_hipMallocPitch_Negative") {
   }
 }
 
+TEST_CASE("Unit_hipMallocPitch_Zero_Dims") {
+  void* ptr = nullptr;
+  size_t pitch = 0;
+
+  SECTION("width == 0") {
+    HIP_CHECK(hipMallocPitch(&ptr, &pitch, 0, 1));
+    REQUIRE(ptr == nullptr);
+  }
+
+  SECTION("height == 0") {
+    HIP_CHECK(hipMallocPitch(&ptr, &pitch, 1, 0));
+    REQUIRE(ptr == nullptr);
+  }
+}
+
 TEST_CASE("Unit_hipMemAllocPitch_Negative") {
   size_t pitch = 0;
   hipDeviceptr_t ptr{};
@@ -366,42 +381,7 @@ static void MemoryAllocDiffSizes(int gpu) {
 static void threadFunc(int gpu) {
   MemoryAllocDiffSizes<float>(gpu);
 }
-/*
- * This testcase verifies the negative scenarios of hipMallocPitch API
- */
-#if 0 //TODO: Review, fix and re-enable test
-TEST_CASE("Unit_hipMallocPitch_Negative") {
-  float* A_d;
-  size_t pitch_A = 0;
-  size_t width{NUM_W * sizeof(float)};
-#if HT_NVIDIA
-  SECTION("NullPtr to Pitched Ptr") {
-    REQUIRE(hipMallocPitch(nullptr,
-            &pitch_A, width, NUM_H) != hipSuccess);
-  }
-
-  SECTION("nullptr to pitch") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           nullptr, width, NUM_H) != hipSuccess);
-  }
-#endif
-  SECTION("Width 0 in hipMallocPitch") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           &pitch_A, 0, NUM_H) == hipSuccess);
-  }
 
-  SECTION("Height 0 in hipMallocPitch") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           &pitch_A, width, 0) == hipSuccess);
-  }
-
-  SECTION("Max int values") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           &pitch_A, std::numeric_limits<int>::max(),
-                           std::numeric_limits<int>::max()) != hipSuccess);
-  }
-}
-#endif
 /*
  * This testcase verifies the basic scenario of
  * hipMallocPitch API for different datatypes
@@ -414,6 +394,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic",
   size_t width{NUM_W * sizeof(TestType)};
   REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
           &pitch_A, width, NUM_H) == hipSuccess);
+  REQUIRE(width <= pitch_A);
   HIP_CHECK(hipFree(A_d));
 }
 

tests/catch/unit/memory/hipPointerGetAttribute.cc

@@ -316,9 +316,8 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") {
         == hipErrorInvalidValue);
   }
   SECTION("Pass invalid attribute") {
-    hipPointer_attribute attr{HIP_POINTER_ATTRIBUTE_DEVICE_POINTER};
-    REQUIRE(hipPointerGetAttribute(&data, attr,
-          reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
+    REQUIRE(hipPointerGetAttribute(&data, static_cast<hipPointer_attribute>(-1),
+                                   reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
   }
 #if HT_AMD
   SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE"