SWDEV-1 - Merge github PRs to amd-staging

Change-Id: I2944a63ddc2eec8dc1403d9790ffffbaec343385


[ROCm/hip-tests commit: 57bc68acb1]

projects/hip-tests/catch/include/cmd_options.hh

@@ -22,6 +22,9 @@ THE SOFTWARE.
 
 #pragma once
 
+#include <cstdint>
+#include <limits>
+
 struct CmdOptions {
   int iterations = 10;
   int warmups = 100;
@@ -29,6 +32,8 @@ struct CmdOptions {
   int cg_iterations = 5;
   bool no_display = false;
   bool progress = false;
+  uint64_t accuracy_iterations = std::numeric_limits<uint32_t>::max() + 1ull;
+  int accuracy_max_memory = 80;
 };
 
 extern CmdOptions cmd_options;

projects/hip-tests/catch/include/hipMallocManagedCommon.hh

@@ -0,0 +1,26 @@
+#include <hip_test_common.hh>
+
+static int HmmAttrPrint() {
+  int managed = 0;
+  INFO(
+      "The following are the attribute values related to HMM for"
+      " device 0:\n");
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
+  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeConcurrentManagedAccess, 0));
+  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
+  HIP_CHECK(
+      hipDeviceGetAttribute(&managed, hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
+  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" << managed);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  if (managed != 1) {
+    WARN(
+        "GPU 0 doesn't support hipDeviceAttributeManagedMemory attribute so defaulting to system "
+        "memory.");
+  }
+  return managed;
+}

projects/hip-tests/catch/include/hip_test_common.hh

@@ -129,6 +129,19 @@ THE SOFTWARE.
     }                                                                                              \
   }
 
+// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError.
+#define HIPRTC_CHECK_ERROR(errorExpr, expectedError)                                               \
+  {                                                                                                \
+    auto localError = errorExpr;                                                                   \
+    INFO("Matching Errors: "                                                                       \
+         << "\n    Expected Error: " << hiprtcGetErrorString(expectedError)                        \
+         << "\n    Expected Code: " << expectedError << '\n'                                       \
+         << "                  Actual Error:   " << hiprtcGetErrorString(localError)               \
+         << "\n    Actual Code:   " << localError << "\nStr: " << #errorExpr                       \
+         << "\n    In File: " << __FILE__ << "\n    At line: " << __LINE__);                       \
+    REQUIRE(localError == expectedError);                                                          \
+  }
+
 #define HIPASSERT(condition)                                                                       \
   if (!(condition)) {                                                                              \
     printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__);                            \
@@ -165,7 +178,7 @@ static inline bool IsGfx11() {
   hipDeviceProp_t props{};
   HIP_CHECK(hipGetDevice(&device));
   HIP_CHECK(hipGetDeviceProperties(&props, device));
-   // Get GCN Arch Name and compare to check if it is gfx11
+  // Get GCN Arch Name and compare to check if it is gfx11
   std::string arch = std::string(props.gcnArchName);
   auto pos = arch.find("gfx11");
   if (pos != std::string::npos)
@@ -173,7 +186,7 @@ static inline bool IsGfx11() {
   else
     return false;
 #else
-  std::cout<<"Have to be either Nvidia or AMD platform, asserting"<<std::endl;
+  std::cout << "Have to be either Nvidia or AMD platform, asserting" << std::endl;
   assert(false);
 #endif
 }
@@ -308,7 +321,7 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB
   launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
                                 std::forward<Args>(packedArgs)...);
 #endif
-HIP_CHECK(hipGetLastError());
+  HIP_CHECK(hipGetLastError());
 }
 
 //---

projects/hip-tests/catch/include/hip_test_defgroups.hh

@@ -39,6 +39,13 @@ THE SOFTWARE.
  * @}
  */
 
+/**
+ * @defgroup AtomicsTest Device Atomics
+ * @{
+ * This section describes tests for the Device Atomic APIs.
+ * @}
+ */
+
 /**
  * @defgroup DeviceLanguageTest Device Language
  * @{
@@ -96,16 +103,23 @@ THE SOFTWARE.
  */
 
 /**
-* @defgroup KernelTest Kernel Functions Management
-* @{
-* This section describes the various kernel functions invocation.
-* @}
-*/
+ * @defgroup KernelTest Kernel Functions Management
+ * @{
+ * This section describes the various kernel functions invocation.
+ * @}
+ */
 
 /**
- * @defgroup AtomicsTest Device Atomics
+ * @defgroup SyncthreadsTest Synchronization Functions
  * @{
- * This section describes tests for the Device Atomic APIs.
+ * This section describes tests for Synchronization Functions.
+ * @}
+ */
+
+/**
+ * @defgroup ThreadfenceTest Memory Fence Functions
+ * @{
+ * This section describes tests for Memory Fence Functions.
  * @}
  */
 
@@ -119,7 +133,8 @@ THE SOFTWARE.
 /**
  * @defgroup PeerToPeerTest PeerToPeer Device Memory Access
  * @{
- * This section describes tests for the PeerToPeer device memory access functions of HIP runtime API.
+ * This section describes tests for the PeerToPeer device memory access functions of HIP runtime
+ * API.
  * @warning PeerToPeer support is experimental.
  * @}
  */
@@ -135,6 +150,7 @@ THE SOFTWARE.
  * @defgroup ShflTest warp shuffle function Management
  * @{
  * This section describes the warp shuffle types & functions of HIP runtime API.
+ * @}
  */
 
 /**
@@ -158,6 +174,13 @@ THE SOFTWARE.
  * @}
  */
 
+/**
+ * @defgroup ModuleTest Module Management
+ * @{
+ * This section describes the module management types & functions of HIP runtime API.
+ * @}
+ */
+
 /**
  * @defgroup TextureTest Texture Management
  * @{
@@ -172,6 +195,13 @@ THE SOFTWARE.
  * @}
  */
 
+ /**
+ * @defgroup MathTest Math Device Functions
+ * @{
+ * This section describes tests for device math functions of HIP runtime API.
+ * @}
+ */
+
 /**
  * @defgroup PrintfTest Printf API Management
  * @{
@@ -192,3 +222,10 @@ THE SOFTWARE.
  * This section describes tests for the Complex type functions.
  * @}
  */
+
+/**
+ * @defgroup VirtualMemoryManagementTest Virtual Memory Management APIs
+ * @{
+ * This section describes the virtual memory management types & functions of HIP runtime API.
+ * @}
+ */

projects/hip-tests/catch/include/memcpy3d_tests_common.hh

@@ -23,7 +23,7 @@ THE SOFTWARE.
 #pragma once
 #pragma clang diagnostic ignored "-Wmissing-field-initializers"
 #pragma clang diagnostic ignored "-Wunused-lambda-capture"
-
+#pragma clang diagnostic ignored "-Wunused-parameter"
 #include <variant>
 
 #include <hip_test_common.hh>
@@ -44,8 +44,9 @@ static inline hipMemcpyKind ReverseMemcpyDirection(const hipMemcpyKind direction
   }
 };
 
-static hipMemcpy3DParms GetMemcpy3DParms(PtrVariant dst_ptr, hipPos dst_pos, PtrVariant src_ptr,
-                                         hipPos src_pos, hipExtent extent, hipMemcpyKind kind) {
+static inline hipMemcpy3DParms GetMemcpy3DParms(PtrVariant dst_ptr, hipPos dst_pos,
+                                                PtrVariant src_ptr, hipPos src_pos,
+                                                hipExtent extent, hipMemcpyKind kind) {
   hipMemcpy3DParms parms = {0};
   if (std::holds_alternative<hipArray_t>(dst_ptr)) {
     parms.dstArray = std::get<hipArray_t>(dst_ptr);
@@ -185,7 +186,7 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, hipStream_t kernel_stream = null
     HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
     if (!can_access_peer) {
       std::string msg = "Skipped as peer access cannot be enabled between devices " +
-                         std::to_string(src_device) + " " + std::to_string(dst_device);
+          std::to_string(src_device) + " " + std::to_string(dst_device);
       HipTest::HIP_SKIP_TEST(msg.c_str());
       return;
     }
@@ -205,7 +206,8 @@ void Memcpy3DDeviceToDeviceShell(F memcpy_func, hipStream_t kernel_stream = null
   // Using dst_alloc width and height to set only the elements that will be copied over to
   // dst_alloc
   Iota<<<blocks, threads_per_block, 0, kernel_stream>>>(src_alloc.ptr(), src_alloc.pitch(),
-                          dst_alloc.width_logical(),dst_alloc.height(), dst_alloc.depth());
+                                                        dst_alloc.width_logical(),
+                                                        dst_alloc.height(), dst_alloc.depth());
   HIP_CHECK(hipGetLastError());
 
   HIP_CHECK(memcpy_func(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
@@ -626,15 +628,14 @@ constexpr auto MemTypeUnified() {
 
 using DrvPtrVariant = std::variant<hipPitchedPtr, hipArray_t>;
 
-template <bool async = false>
-hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr,
-                              hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
-                              hipStream_t stream = nullptr) {
+static inline HIP_MEMCPY3D GetDrvMemcpy3DParms(DrvPtrVariant dst_ptr, hipPos dst_pos,
+                                               DrvPtrVariant src_ptr, hipPos src_pos,
+                                               hipExtent extent, hipMemcpyKind kind) {
   HIP_MEMCPY3D parms = {0};
 
   if (std::holds_alternative<hipArray_t>(dst_ptr)) {
     parms.dstMemoryType = hipMemoryTypeArray;
-    parms.dstArray = std::get<hipArray_t>(dst_ptr);  
+    parms.dstArray = std::get<hipArray_t>(dst_ptr);
   } else {
     auto ptr = std::get<hipPitchedPtr>(dst_ptr);
     parms.dstPitch = ptr.pitch;
@@ -694,6 +695,84 @@ hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVaria
   parms.dstY = dst_pos.y;
   parms.dstZ = dst_pos.z;
 
+  return parms;
+}
+
+static inline bool operator==(const HIP_MEMCPY3D& lhs, const HIP_MEMCPY3D& rhs) {
+  bool pos_eq = lhs.dstXInBytes == rhs.dstXInBytes && lhs.dstY == rhs.dstY &&
+      lhs.dstZ == rhs.dstZ && lhs.srcXInBytes == rhs.srcXInBytes && lhs.srcY == rhs.srcY &&
+      lhs.srcZ == rhs.srcZ;
+  bool extent_eq =
+      lhs.WidthInBytes == rhs.WidthInBytes && lhs.Height == rhs.Height && lhs.Depth == rhs.Depth;
+  bool mem_eq = true;
+  if (lhs.dstArray) {
+    mem_eq = lhs.dstArray == rhs.dstArray && lhs.dstMemoryType == rhs.dstMemoryType;
+  } else {
+    mem_eq = lhs.dstPitch == rhs.dstPitch && lhs.dstMemoryType == rhs.dstMemoryType;
+  }
+  if (lhs.srcArray) {
+    mem_eq = lhs.srcArray == rhs.srcArray && lhs.srcMemoryType == rhs.srcMemoryType;
+  } else {
+    mem_eq = lhs.srcPitch == rhs.srcPitch && lhs.srcMemoryType == rhs.srcMemoryType;
+  }
+  if (lhs.dstDevice) {
+    mem_eq = mem_eq && (lhs.dstDevice == rhs.dstDevice);
+  }
+  if (lhs.dstHost) {
+    mem_eq = mem_eq && (lhs.dstDevice == rhs.dstDevice);
+  }
+  if (lhs.srcDevice) {
+    mem_eq = mem_eq && (lhs.srcDevice == rhs.srcDevice);
+  }
+  if (lhs.srcHost) {
+    mem_eq = mem_eq && (lhs.srcHost == rhs.srcHost);
+  }
+
+  return pos_eq && extent_eq && mem_eq;
+}
+
+// APIs hipDrvGraphMemcpyNodeGetParams, hipDrvGraphMemcpyNodeSetParams are yet to be implemented in HIP runtime.
+#if 0
+template <bool set_params = false>
+hipError_t DrvMemcpy3DGraphWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr,
+                                   hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
+                                   hipCtx_t context, hipStream_t stream = nullptr) {
+  auto parms = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+
+  hipGraph_t g = nullptr;
+  HIP_CHECK(hipGraphCreate(&g, 0));
+  hipGraphNode_t node = nullptr;
+  if constexpr (set_params) {
+    auto reversed_parms = GetDrvMemcpy3DParms(src_ptr, src_pos, dst_ptr, dst_pos, extent,
+                                              ReverseMemcpyDirection(kind));
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, g, nullptr, 0, &reversed_parms, context));
+    HIP_CHECK(hipDrvGraphMemcpyNodeSetParams(node, &parms));
+  } else {
+    HIP_CHECK(hipDrvGraphAddMemcpyNode(&node, g, nullptr, 0, &parms, context));
+  }
+
+  HIP_MEMCPY3D retrieved_params = {0};
+  HIP_CHECK(hipDrvGraphMemcpyNodeGetParams(node, &retrieved_params));
+  REQUIRE(parms == retrieved_params);
+
+  hipGraphExec_t graph_exec = nullptr;
+  HIP_CHECK(hipGraphInstantiate(&graph_exec, g, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread));
+  HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
+
+  HIP_CHECK(hipGraphExecDestroy(graph_exec));
+  HIP_CHECK(hipGraphDestroy(g));
+
+  return hipSuccess;
+}
+#endif //if 0
+
+template <bool async = false>
+hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr,
+                              hipPos src_pos, hipExtent extent, hipMemcpyKind kind,
+                              hipStream_t stream = nullptr) {
+  auto parms = GetDrvMemcpy3DParms(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind);
+
   if constexpr (async) {
     return hipDrvMemcpy3DAsync(&parms, stream);
   } else {
@@ -805,4 +884,4 @@ void DrvMemcpy3DArrayDeviceShell(F memcpy_func, const hipStream_t kernel_stream
   };
   PitchedMemoryVerify(host_alloc.ptr(), extent.width, extent.width / sizeof(int), extent.height,
                       extent.depth, f);
-}
+}

projects/hip-tests/catch/include/resource_guards.hh

@@ -35,15 +35,15 @@ enum class LinearAllocs {
 inline std::string to_string(const LinearAllocs allocation_type) {
   switch (allocation_type) {
     case LinearAllocs::malloc:
-      return "host pageable";
+      return "malloc";
     case LinearAllocs::mallocAndRegister:
-      return "registered";
+      return "malloc + hipHostRegister";
     case LinearAllocs::hipHostMalloc:
-      return "host pinned";
+      return "hipHostMalloc";
     case LinearAllocs::hipMalloc:
-      return "device malloc";
+      return "hipMalloc";
     case LinearAllocs::hipMallocManaged:
-      return "managed";
+      return "hipMallocManaged";
     default:
       return "unknown alloc type";
   }
@@ -83,24 +83,38 @@ template <typename T> class LinearAllocGuard {
 
   LinearAllocGuard(const LinearAllocGuard&) = delete;
 
-  LinearAllocGuard(LinearAllocGuard&& o)
-      : allocation_type_{o.allocation_type_}, ptr_{o.ptr_}, host_ptr_{o.host_ptr_} {
-    o.allocation_type_ = LinearAllocs::noAlloc;
-    o.ptr_ = nullptr;
-    o.host_ptr_ = nullptr;
-  }
+  LinearAllocGuard(LinearAllocGuard&& o) { *this = std::move(o); }
 
   LinearAllocGuard& operator=(LinearAllocGuard&& o) {
-    allocation_type_ = o.allocation_type_;
-    ptr_ = o.ptr_;
-    host_ptr_ = o.host_ptr_;
+    if (this != &o) {
+      dealloc();
 
-    o.allocation_type_ = LinearAllocs::noAlloc;
-    o.ptr_ = nullptr;
-    o.host_ptr_ = nullptr;
+      allocation_type_ = o.allocation_type_;
+      ptr_ = o.ptr_;
+      host_ptr_ = o.host_ptr_;
+
+      o.allocation_type_ = LinearAllocs::noAlloc;
+      o.ptr_ = nullptr;
+      o.host_ptr_ = nullptr;
+    }
+
+    return *this;
   }
 
-  ~LinearAllocGuard() {
+  ~LinearAllocGuard() { dealloc(); }
+
+  T* ptr() const { return ptr_; };
+  T* host_ptr() const { return host_ptr_; }
+
+ private:
+  LinearAllocs allocation_type_ = LinearAllocs::noAlloc;
+  T* ptr_ = nullptr;
+  T* host_ptr_ = nullptr;
+
+  void dealloc() {
+    if (ptr_ == nullptr) {
+      return;
+    }
     // No Catch macros, don't want to possibly throw in the destructor
     if (ptr_ != nullptr) {
       switch (allocation_type_) {
@@ -123,14 +137,6 @@ template <typename T> class LinearAllocGuard {
       }
     }
   }
-
-  T* ptr() const { return ptr_; };
-  T* host_ptr() const { return host_ptr_; }
-
- private:
-  LinearAllocs allocation_type_ = LinearAllocs::noAlloc;
-  T* ptr_ = nullptr;
-  T* host_ptr_ = nullptr;
 };
 
 template <typename T> class LinearAllocGuardMultiDim {
@@ -210,6 +216,42 @@ template <typename T> class ArrayAllocGuard {
   const hipExtent extent_;
 };
 
+template <typename T> class MipmappedArrayAllocGuard {
+ public:
+  // extent should contain logical width
+  MipmappedArrayAllocGuard(const hipExtent extent, const unsigned int levels,
+                           const unsigned int flags)
+      : extent_{extent}, levels_{levels} {
+    hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
+    HIP_CHECK(hipMallocMipmappedArray(&ptr_, &desc, extent_, levels_, flags));
+  }
+
+  MipmappedArrayAllocGuard(const hipExtent extent, const unsigned int flags = 0u)
+      : MipmappedArrayAllocGuard{extent, 1, flags} {}
+
+  ~MipmappedArrayAllocGuard() { static_cast<void>(hipFreeMipmappedArray(ptr_)); }
+
+  MipmappedArrayAllocGuard(const MipmappedArrayAllocGuard&) = delete;
+  MipmappedArrayAllocGuard(MipmappedArrayAllocGuard&&) = delete;
+
+  hipMipmappedArray_t ptr() const { return ptr_; }
+
+  hipArray_t GetLevel(unsigned int level) {
+    hipArray_t ret;
+    HIP_CHECK(hipGetMipmappedArrayLevel(&ret, ptr_, level));
+    return ret;
+  }
+
+  hipExtent extent() const { return extent_; }
+
+  unsigned int levels() const { return levels_; }
+
+ private:
+  hipMipmappedArray_t ptr_ = nullptr;
+  const hipExtent extent_;
+  const unsigned int levels_;
+};
+
 template <typename T> class DrvArrayAllocGuard {
  public:
   // extent should contain width in bytes
@@ -266,24 +308,24 @@ class StreamGuard {
 
   StreamGuard(const StreamGuard&) = delete;
 
-  StreamGuard(StreamGuard&& o)
-      : stream_type_{o.stream_type_}, flags_{o.flags_}, priority_{o.priority_}, stream_{o.stream_} {
-    o.stream_type_ = Streams::nullstream;
-    o.flags_ = 0u;
-    o.priority_ = 0;
-    o.stream_ = nullptr;
-  }
+  StreamGuard(StreamGuard&& o) { *this = std::move(o); }
 
   StreamGuard& operator=(StreamGuard&& o) {
-    stream_type_ = o.stream_type_;
-    flags_ = o.flags_;
-    priority_ = o.priority_;
-    stream_ = o.stream_;
+    if (this != &o) {
+      if (stream_type_ == Streams::created) {
+        static_cast<void>(hipStreamDestroy(stream_));
+      }
 
-    o.stream_type_ = Streams::nullstream;
-    o.flags_ = 0u;
-    o.priority_ = 0;
-    o.stream_ = nullptr;
+      stream_type_ = o.stream_type_;
+      flags_ = o.flags_;
+      priority_ = o.priority_;
+      stream_ = o.stream_;
+
+      o.stream_type_ = Streams::nullstream;
+      o.flags_ = 0u;
+      o.priority_ = 0;
+      o.stream_ = nullptr;
+    }
 
     return *this;
   }

projects/hip-tests/catch/include/utils.hh

@@ -170,7 +170,7 @@ inline bool DeviceAttributesSupport(const int device, Attributes... attributes)
   return (... && DeviceAttributeSupport(device, attributes));
 }
 
-inline int GetDeviceAttribute(int device, const hipDeviceAttribute_t attr) {
+inline int GetDeviceAttribute(const hipDeviceAttribute_t attr, int device) {
   int value = 0;
   HIP_CHECK(hipDeviceGetAttribute(&value, attr, device));
   return value;