Merge branch 'main' into fix-coverity/sanitizer-layer

Author: AllanZyne

include/ur_api.h

@@ -3237,13 +3237,16 @@ typedef enum ur_usm_advice_flag_t {
     UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_HOST = UR_BIT(12),        ///< Removes the affect of ::UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_HOST
     UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION_HOST = UR_BIT(13),   ///< Hint that the preferred memory location is the host
     UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST = UR_BIT(14), ///< Removes the affect of ::UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION_HOST
+    UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY = UR_BIT(15),       ///< Hint that memory coherence will be coarse-grained (up-to-date only at
+                                                                   ///< kernel boundaries)
+    UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY = UR_BIT(16),     ///< Removes the effect of ::UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY
     /// @cond
     UR_USM_ADVICE_FLAG_FORCE_UINT32 = 0x7fffffff
     /// @endcond
 
 } ur_usm_advice_flag_t;
 /// @brief Bit Mask for validating ur_usm_advice_flags_t
-#define UR_USM_ADVICE_FLAGS_MASK 0xffff8000
+#define UR_USM_ADVICE_FLAGS_MASK 0xfffe0000
 
 ///////////////////////////////////////////////////////////////////////////////
 /// @brief Handle of USM pool

include/ur_print.hpp

@@ -6277,6 +6277,12 @@ inline std::ostream &operator<<(std::ostream &os, ur_usm_advice_flag_t value) {
     case UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST:
         os << "UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST";
         break;
+    case UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY:
+        os << "UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY";
+        break;
+    case UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY:
+        os << "UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY";
+        break;
     default:
         os << "unknown enumerator";
         break;
@@ -6441,6 +6447,26 @@ inline ur_result_t printFlag<ur_usm_advice_flag_t>(std::ostream &os, uint32_t fl
         }
         os << UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST;
     }
+
+    if ((val & UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY) == (uint32_t)UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY) {
+        val ^= (uint32_t)UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY;
+        if (!first) {
+            os << " | ";
+        } else {
+            first = false;
+        }
+        os << UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY;
+    }
+
+    if ((val & UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY) == (uint32_t)UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY) {
+        val ^= (uint32_t)UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY;
+        if (!first) {
+            os << " | ";
+        } else {
+            first = false;
+        }
+        os << UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY;
+    }
     if (val != 0) {
         std::bitset<32> bits(val);
         if (!first) {

scripts/core/usm.yml

@@ -126,6 +126,12 @@ etors:
     - name: CLEAR_PREFERRED_LOCATION_HOST
       value: "$X_BIT(14)"
       desc: "Removes the affect of $X_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION_HOST"
+    - name: SET_NON_COHERENT_MEMORY
+      value: "$X_BIT(15)"
+      desc: "Hint that memory coherence will be coarse-grained (up-to-date only at kernel boundaries)"
+    - name: CLEAR_NON_COHERENT_MEMORY
+      value: "$X_BIT(16)"
+      desc: "Removes the effect of $X_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY"
 --- #--------------------------------------------------------------------------
 type: handle
 desc: "Handle of USM pool"

source/adapters/cuda/enqueue.cpp

@@ -114,10 +114,13 @@ ur_result_t setCuMemAdvise(CUdeviceptr DevPtr, size_t Size,
     }
   }
 
-  std::array<ur_usm_advice_flags_t, 4> UnmappedMemAdviceFlags = {
+  std::array<ur_usm_advice_flags_t, 6> UnmappedMemAdviceFlags = {
       UR_USM_ADVICE_FLAG_SET_NON_ATOMIC_MOSTLY,
       UR_USM_ADVICE_FLAG_CLEAR_NON_ATOMIC_MOSTLY,
-      UR_USM_ADVICE_FLAG_BIAS_CACHED, UR_USM_ADVICE_FLAG_BIAS_UNCACHED};
+      UR_USM_ADVICE_FLAG_BIAS_CACHED,
+      UR_USM_ADVICE_FLAG_BIAS_UNCACHED,
+      UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY,
+      UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY};
 
   for (auto &UnmappedFlag : UnmappedMemAdviceFlags) {
     if (URAdviceFlags & UnmappedFlag) {

source/adapters/hip/device.hpp

@@ -26,12 +26,37 @@ struct ur_device_handle_t_ {
   ur_platform_handle_t Platform;
   hipCtx_t HIPContext;
   uint32_t DeviceIndex;
+  int MaxWorkGroupSize{0};
+  int MaxBlockDimX{0};
+  int MaxBlockDimY{0};
+  int MaxBlockDimZ{0};
+  int DeviceMaxLocalMem{0};
+  int ManagedMemSupport{0};
+  int ConcurrentManagedAccess{0};
 
 public:
   ur_device_handle_t_(native_type HipDevice, hipCtx_t Context,
                       ur_platform_handle_t Platform, uint32_t DeviceIndex)
       : HIPDevice(HipDevice), RefCount{1}, Platform(Platform),
-        HIPContext(Context), DeviceIndex(DeviceIndex) {}
+        HIPContext(Context), DeviceIndex(DeviceIndex) {
+
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &MaxWorkGroupSize, hipDeviceAttributeMaxThreadsPerBlock, HIPDevice));
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &MaxBlockDimX, hipDeviceAttributeMaxBlockDimX, HIPDevice));
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &MaxBlockDimY, hipDeviceAttributeMaxBlockDimY, HIPDevice));
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &MaxBlockDimZ, hipDeviceAttributeMaxBlockDimZ, HIPDevice));
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &DeviceMaxLocalMem, hipDeviceAttributeMaxSharedMemoryPerBlock,
+        HIPDevice));
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &ManagedMemSupport, hipDeviceAttributeManagedMemory, HIPDevice));
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &ConcurrentManagedAccess, hipDeviceAttributeConcurrentManagedAccess,
+        HIPDevice));
+  }
 
   ~ur_device_handle_t_() noexcept(false) {
     UR_CHECK_ERROR(hipDevicePrimaryCtxRelease(HIPDevice));
@@ -48,6 +73,22 @@ struct ur_device_handle_t_ {
   // Returns the index of the device relative to the other devices in the same
   // platform
   uint32_t getIndex() const noexcept { return DeviceIndex; };
+
+  int getMaxWorkGroupSize() const noexcept { return MaxWorkGroupSize; };
+
+  int getMaxBlockDimX() const noexcept { return MaxBlockDimX; };
+
+  int getMaxBlockDimY() const noexcept { return MaxBlockDimY; };
+
+  int getMaxBlockDimZ() const noexcept { return MaxBlockDimZ; };
+
+  int getDeviceMaxLocalMem() const noexcept { return DeviceMaxLocalMem; };
+
+  int getManagedMemSupport() const noexcept { return ManagedMemSupport; };
+
+  int getConcurrentManagedAccess() const noexcept {
+    return ConcurrentManagedAccess;
+  };
 };
 
 int getAttribute(ur_device_handle_t Device, hipDeviceAttribute_t Attribute);

source/adapters/hip/enqueue.cpp

@@ -75,30 +75,46 @@ ur_result_t setHipMemAdvise(const void *DevPtr, const size_t Size,
   if (URAdviceFlags &
       (UR_USM_ADVICE_FLAG_SET_NON_ATOMIC_MOSTLY |
        UR_USM_ADVICE_FLAG_CLEAR_NON_ATOMIC_MOSTLY |
-       UR_USM_ADVICE_FLAG_BIAS_CACHED | UR_USM_ADVICE_FLAG_BIAS_UNCACHED)) {
+       UR_USM_ADVICE_FLAG_BIAS_CACHED | UR_USM_ADVICE_FLAG_BIAS_UNCACHED
+#if !defined(__HIP_PLATFORM_AMD__)
+       | UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY |
+       UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY
+#endif
+       )) {
     return UR_RESULT_ERROR_INVALID_ENUMERATION;
   }
 
   using ur_to_hip_advice_t = std::pair<ur_usm_advice_flags_t, hipMemoryAdvise>;
 
-  static constexpr std::array<ur_to_hip_advice_t, 6>
-      URToHIPMemAdviseDeviceFlags{
-          std::make_pair(UR_USM_ADVICE_FLAG_SET_READ_MOSTLY,
-                         hipMemAdviseSetReadMostly),
-          std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_READ_MOSTLY,
-                         hipMemAdviseUnsetReadMostly),
-          std::make_pair(UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION,
-                         hipMemAdviseSetPreferredLocation),
-          std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION,
-                         hipMemAdviseUnsetPreferredLocation),
-          std::make_pair(UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE,
-                         hipMemAdviseSetAccessedBy),
-          std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE,
-                         hipMemAdviseUnsetAccessedBy),
-      };
-  for (auto &FlagPair : URToHIPMemAdviseDeviceFlags) {
-    if (URAdviceFlags & FlagPair.first) {
-      UR_CHECK_ERROR(hipMemAdvise(DevPtr, Size, FlagPair.second, Device));
+#if defined(__HIP_PLATFORM_AMD__)
+  constexpr size_t DeviceFlagCount = 8;
+#else
+  constexpr size_t DeviceFlagCount = 6;
+#endif
+  static constexpr std::array<ur_to_hip_advice_t, DeviceFlagCount>
+      URToHIPMemAdviseDeviceFlags {
+    std::make_pair(UR_USM_ADVICE_FLAG_SET_READ_MOSTLY,
+                   hipMemAdviseSetReadMostly),
+        std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_READ_MOSTLY,
+                       hipMemAdviseUnsetReadMostly),
+        std::make_pair(UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION,
+                       hipMemAdviseSetPreferredLocation),
+        std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION,
+                       hipMemAdviseUnsetPreferredLocation),
+        std::make_pair(UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE,
+                       hipMemAdviseSetAccessedBy),
+        std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE,
+                       hipMemAdviseUnsetAccessedBy),
+#if defined(__HIP_PLATFORM_AMD__)
+        std::make_pair(UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY,
+                       hipMemAdviseSetCoarseGrain),
+        std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_NON_COHERENT_MEMORY,
+                       hipMemAdviseUnsetCoarseGrain),
+#endif
+  };
+  for (const auto &[URAdvice, HIPAdvice] : URToHIPMemAdviseDeviceFlags) {
+    if (URAdviceFlags & URAdvice) {
+      UR_CHECK_ERROR(hipMemAdvise(DevPtr, Size, HIPAdvice, Device));
     }
   }
 
@@ -113,10 +129,9 @@ ur_result_t setHipMemAdvise(const void *DevPtr, const size_t Size,
                      hipMemAdviseUnsetAccessedBy),
   };
 
-  for (auto &FlagPair : URToHIPMemAdviseHostFlags) {
-    if (URAdviceFlags & FlagPair.first) {
-      UR_CHECK_ERROR(
-          hipMemAdvise(DevPtr, Size, FlagPair.second, hipCpuDeviceId));
+  for (const auto &[URAdvice, HIPAdvice] : URToHIPMemAdviseHostFlags) {
+    if (URAdviceFlags & URAdvice) {
+      UR_CHECK_ERROR(hipMemAdvise(DevPtr, Size, HIPAdvice, hipCpuDeviceId));
     }
   }
 
@@ -300,15 +315,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   bool ProvidedLocalWorkGroupSize = (pLocalWorkSize != nullptr);
 
   {
-    ur_result_t Result = urDeviceGetInfo(
-        hQueue->Device, UR_DEVICE_INFO_MAX_WORK_ITEM_SIZES,
-        sizeof(MaxThreadsPerBlock), MaxThreadsPerBlock, nullptr);
-    UR_ASSERT(Result == UR_RESULT_SUCCESS, Result);
+    MaxThreadsPerBlock[0] = hQueue->Device->getMaxBlockDimX();
+    MaxThreadsPerBlock[1] = hQueue->Device->getMaxBlockDimY();
+    MaxThreadsPerBlock[2] = hQueue->Device->getMaxBlockDimZ();
 
-    Result =
-        urDeviceGetInfo(hQueue->Device, UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE,
-                        sizeof(MaxWorkGroupSize), &MaxWorkGroupSize, nullptr);
-    UR_ASSERT(Result == UR_RESULT_SUCCESS, Result);
+    MaxWorkGroupSize = hQueue->Device->getMaxWorkGroupSize();
 
     // The MaxWorkGroupSize = 1024 for AMD GPU
     // The MaxThreadsPerBlock = {1024, 1024, 1024}
@@ -423,11 +434,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
                         : (LocalMemSzPtrPI ? LocalMemSzPtrPI : nullptr);
 
     if (LocalMemSzPtr) {
-      int DeviceMaxLocalMem = 0;
-      UR_CHECK_ERROR(hipDeviceGetAttribute(
-          &DeviceMaxLocalMem, hipDeviceAttributeMaxSharedMemoryPerBlock,
-          Dev->get()));
-
+      int DeviceMaxLocalMem = Dev->getDeviceMaxLocalMem();
       static const int EnvVal = std::atoi(LocalMemSzPtr);
       if (EnvVal <= 0 || EnvVal > DeviceMaxLocalMem) {
         setErrorMessage(LocalMemSzPtrUR ? "Invalid value specified for "
@@ -1484,7 +1491,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(
 
     // If the device does not support managed memory access, we can't set
     // mem_advise.
-    if (!getAttribute(Device, hipDeviceAttributeManagedMemory)) {
+    if (!Device->getManagedMemSupport()) {
       releaseEvent();
       setErrorMessage("mem_advise ignored as device does not support "
                       "managed memory access",
@@ -1558,7 +1565,7 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size,
 
     // If the device does not support managed memory access, we can't set
     // mem_advise.
-    if (!getAttribute(Device, hipDeviceAttributeManagedMemory)) {
+    if (!Device->getManagedMemSupport()) {
       releaseEvent();
       setErrorMessage("mem_advise ignored as device does not support "
                       "managed memory access",
@@ -1575,7 +1582,7 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size,
                   UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE |
                   UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE |
                   UR_USM_ADVICE_FLAG_DEFAULT)) {
-      if (!getAttribute(Device, hipDeviceAttributeConcurrentManagedAccess)) {
+      if (!Device->getConcurrentManagedAccess()) {
         releaseEvent();
         setErrorMessage("mem_advise ignored as device does not support "
                         "concurrent managed access",
@@ -1598,6 +1605,10 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size,
           pMem, size, hipMemAdviseUnsetPreferredLocation, DeviceID));
       UR_CHECK_ERROR(
           hipMemAdvise(pMem, size, hipMemAdviseUnsetAccessedBy, DeviceID));
+#if defined(__HIP_PLATFORM_AMD__)
+      UR_CHECK_ERROR(
+          hipMemAdvise(pMem, size, hipMemAdviseUnsetCoarseGrain, DeviceID));
+#endif
     } else {
       Result = setHipMemAdvise(HIPDevicePtr, size, advice, DeviceID);
       // UR_RESULT_ERROR_INVALID_ENUMERATION is returned when using a valid but
@@ -1663,8 +1674,57 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy2D(
       UR_CHECK_ERROR(RetImplEvent->start());
     }
 
+    // There is an issue with hipMemcpy2D* when hipMemcpyDefault is used, which
+    // makes the HIP runtime not correctly derive the copy kind (direction) for
+    // the copies since ROCm 5.6.0+. See: https://github.com/ROCm/clr/issues/40
+    // TODO: Add maximum HIP_VERSION when bug has been fixed.
+#if HIP_VERSION >= 50600000
+    hipPointerAttribute_t srcAttribs{};
+    hipPointerAttribute_t dstAttribs{};
+
+    bool srcIsSystemAlloc{false};
+    bool dstIsSystemAlloc{false};
+
+    hipError_t hipRes{};
+    // hipErrorInvalidValue returned from hipPointerGetAttributes for a non-null
+    // pointer refers to an OS-allocation, hence pageable host memory. However,
+    // this means we cannot rely on the attributes result, hence we mark system
+    // pageable memory allocation manually as host memory. The HIP runtime can
+    // handle the registering/unregistering of the memory as long as the right
+    // copy-kind (direction) is provided to hipMemcpy2DAsync for this case.
+    hipRes = hipPointerGetAttributes(&srcAttribs, (const void *)pSrc);
+    if (hipRes == hipErrorInvalidValue && pSrc)
+      srcIsSystemAlloc = true;
+    hipRes = hipPointerGetAttributes(&dstAttribs, (const void *)pDst);
+    if (hipRes == hipErrorInvalidValue && pDst)
+      dstIsSystemAlloc = true;
+
+    const unsigned int srcMemType{srcAttribs.type};
+    const unsigned int dstMemType{dstAttribs.type};
+
+    const bool srcIsHost{(srcMemType == hipMemoryTypeHost) || srcIsSystemAlloc};
+    const bool srcIsDevice{srcMemType == hipMemoryTypeDevice};
+    const bool dstIsHost{(dstMemType == hipMemoryTypeHost) || dstIsSystemAlloc};
+    const bool dstIsDevice{dstMemType == hipMemoryTypeDevice};
+
+    unsigned int cpyKind{};
+    if (srcIsHost && dstIsHost)
+      cpyKind = hipMemcpyHostToHost;
+    else if (srcIsHost && dstIsDevice)
+      cpyKind = hipMemcpyHostToDevice;
+    else if (srcIsDevice && dstIsHost)
+      cpyKind = hipMemcpyDeviceToHost;
+    else if (srcIsDevice && dstIsDevice)
+      cpyKind = hipMemcpyDeviceToDevice;
+    else
+      cpyKind = hipMemcpyDefault;
+
+    UR_CHECK_ERROR(hipMemcpy2DAsync(pDst, dstPitch, pSrc, srcPitch, width,
+                                    height, (hipMemcpyKind)cpyKind, HIPStream));
+#else
     UR_CHECK_ERROR(hipMemcpy2DAsync(pDst, dstPitch, pSrc, srcPitch, width,
                                     height, hipMemcpyDefault, HIPStream));
+#endif
 
     if (phEvent) {
       UR_CHECK_ERROR(RetImplEvent->record());

source/adapters/hip/usm.cpp

@@ -78,9 +78,10 @@ USMFreeImpl([[maybe_unused]] ur_context_handle_t hContext, void *pMem) {
 #else
     const auto Type = hipPointerAttributeType.memoryType;
 #endif
-    UR_ASSERT(Type == hipMemoryTypeDevice || Type == hipMemoryTypeHost,
+    UR_ASSERT(Type == hipMemoryTypeDevice || Type == hipMemoryTypeHost ||
+                  Type == hipMemoryTypeManaged,
               UR_RESULT_ERROR_INVALID_MEM_OBJECT);
-    if (Type == hipMemoryTypeDevice) {
+    if (Type == hipMemoryTypeDevice || Type == hipMemoryTypeManaged) {
       UR_CHECK_ERROR(hipFree(pMem));
     }
     if (Type == hipMemoryTypeHost) {

test/conformance/enqueue/urEnqueueUSMAdvise.cpp

@@ -68,3 +68,11 @@ TEST_P(urEnqueueUSMAdviseTest, InvalidSizeTooLarge) {
                      urEnqueueUSMAdvise(queue, ptr, allocation_size * 2,
                                         UR_USM_ADVICE_FLAG_DEFAULT, nullptr));
 }
+
+TEST_P(urEnqueueUSMAdviseTest, NonCoherentDeviceMemorySuccessOrWarning) {
+    ur_result_t result =
+        urEnqueueUSMAdvise(queue, ptr, allocation_size,
+                           UR_USM_ADVICE_FLAG_SET_NON_COHERENT_MEMORY, nullptr);
+    ASSERT_EQ(result,
+              result & (UR_RESULT_SUCCESS | UR_RESULT_ERROR_ADAPTER_SPECIFIC));
+}