EXSWHTEC-83 - Implement new and reimplement existing tests for hipMemAdvise

tests/catch/hipTestMain/config/config_amd_linux_common.json

@@ -10,6 +10,8 @@
 	   "Unit_hipGetDeviceFlags_Positive_Context",
 	   "Unit_hipIpcCloseMemHandle_Negative_Close_In_Originating_Process",
 	   "Unit_hipIpcOpenMemHandle_Negative_Open_In_Creating_Process",
-	   "Unit_hipDeviceGetPCIBusId_Negative_PartialFill"
+	   "Unit_hipDeviceGetPCIBusId_Negative_PartialFill",
+	   "Unit_hipMemAdvise_AccessedBy_All_Devices",
+	   "Unit_hipMemAdvise_No_Flag_Interference"
 	]
 }

tests/catch/hipTestMain/config/config_amd_windows_common.json

@@ -79,6 +79,8 @@
 	   "Unit_hipIpcOpenMemHandle_Negative_Open_In_Creating_Process",
 	   "Unit_hipDeviceGetPCIBusId_Negative_PartialFill",
 	   "Unit_hipDeviceGetSharedMemConfig_Positive_Basic",
-	   "Unit_hipDeviceGetSharedMemConfig_Positive_Threaded"
+	   "Unit_hipDeviceGetSharedMemConfig_Positive_Threaded",
+	   "Unit_hipMemAdvise_AccessedBy_All_Devices",
+	   "Unit_hipMemAdvise_No_Flag_Interference"
 	]
 }

tests/catch/include/hip_array_common.hh

@@ -0,0 +1,84 @@
+/*
+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>
+
+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> {};

tests/catch/include/resource_guards.hh

@@ -19,6 +19,7 @@ THE SOFTWARE.
 
 #pragma once
 
+#include <hip_array_common.hh>
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
 
@@ -80,17 +81,121 @@ 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; }
+};
+
+template <typename T> class ArrayAllocGuard {
+ public:
+  // extent should contain logical width
+  ArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} {
+    hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
+    HIP_CHECK(hipMalloc3DArray(&ptr_, &desc, extent_, flags));
+  }
+
+  ~ArrayAllocGuard() { static_cast<void>(hipFreeArray(ptr_)); }
+
+  ArrayAllocGuard(const ArrayAllocGuard&) = delete;
+  ArrayAllocGuard(ArrayAllocGuard&&) = delete;
+
+  hipArray_t ptr() const { return ptr_; }
+
+  hipExtent extent() const { return extent_; }
+
+ private:
+  hipArray_t ptr_ = nullptr;
+  const hipExtent extent_;
+};
+
+template <typename T> class DrvArrayAllocGuard {
+ public:
+  // extent should contain width in bytes
+  DrvArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u) : extent_{extent} {
+    HIP_ARRAY3D_DESCRIPTOR desc{};
+    using vec_info = vector_info<T>;
+    desc.Format = vec_info::format;
+    desc.NumChannels = vec_info::size;
+    desc.Width = extent_.width / sizeof(T);
+    desc.Height = extent_.height;
+    desc.Depth = extent_.depth;
+    desc.Flags = flags;
+    HIP_CHECK(hipArray3DCreate(&ptr_, &desc));
+  }
+
+  ~DrvArrayAllocGuard() { static_cast<void>(hipArrayDestroy(ptr_)); }
+
+  DrvArrayAllocGuard(const DrvArrayAllocGuard&) = delete;
+  DrvArrayAllocGuard(DrvArrayAllocGuard&&) = delete;
+
+  hiparray ptr() const { return ptr_; }
+
+  hipExtent extent() const { return extent_; }
+
+ private:
+  hiparray ptr_ = nullptr;
+  const hipExtent extent_;
+};
+
 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 (size_t z = 0; z < depth; ++z) {
+    for (size_t y = 0; y < height; ++y) {
+      for (size_t 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 (size_t z = 0; z < depth; ++z) {
+    for (size_t y = 0; y < height; ++y) {
+      for (size_t 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/CMakeLists.txt

@@ -97,6 +97,7 @@ set(TEST_SRC
     hipMemsetSync.cc
     hipMemsetAsync.cc
     hipMemAdvise.cc
+    hipMemAdvise_old.cc
     hipMemRangeGetAttributes.cc
 )
 else()
@@ -171,6 +172,7 @@ set(TEST_SRC
     hipMemsetSync.cc
     hipMemsetAsync.cc
     hipMemAdvise.cc
+    hipMemAdvise_old.cc
     hipMemRangeGetAttributes.cc
 )
 endif()

tests/catch/unit/memory/hipArray3DCreate.cc

@@ -20,6 +20,7 @@ THE SOFTWARE.
 #include <limits>
 #include "DriverContext.hh"
 #include "hipArrayCommon.hh"
+#include "hip_array_common.hh"
 #include "hip_test_common.hh"
 
 namespace {

tests/catch/unit/memory/hipArrayCommon.hh

@@ -26,66 +26,6 @@ THE SOFTWARE.
 
 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>

tests/catch/unit/memory/hipArrayCreate.cc

@@ -27,6 +27,7 @@ hipArrayCreate API test scenarios
 #include <array>
 #include <numeric>
 #include <hip_test_common.hh>
+#include <hip_array_common.hh>
 #include "hipArrayCommon.hh"
 #include "DriverContext.hh"