Merge branch 'main' into deprecated-header

Author: Hugh Delaney

source/adapters/hip/common.hpp

@@ -204,3 +204,38 @@ template <typename T> class ReleaseGuard {
   /// UR object.
   void dismiss() { Captive = nullptr; }
 };
+
+// Helper method to return a (non-null) pointer's attributes, or std::nullopt in
+// the case that the pointer is unknown to the HIP subsystem.
+inline static std::optional<hipPointerAttribute_t>
+getPointerAttributes(const void *pMem) {
+  // do not throw if hipPointerGetAttributes returns hipErrorInvalidValue
+  hipPointerAttribute_t hipPointerAttributes;
+  hipError_t Ret = hipPointerGetAttributes(&hipPointerAttributes, pMem);
+  if (Ret == hipErrorInvalidValue && pMem) {
+    // pointer non-null but not known to the HIP subsystem
+    return std::nullopt;
+  }
+  // Direct usage of the function, instead of UR_CHECK_ERROR, so we can get
+  // the line offset.
+  checkErrorUR(Ret, __func__, __LINE__ - 7, __FILE__);
+  // ROCm 6.0.0 introduces hipMemoryTypeUnregistered in the hipMemoryType
+  // enum to mark unregistered allocations (i.e., via system allocators).
+#if HIP_VERSION_MAJOR >= 6
+  if (hipPointerAttributes.type == hipMemoryTypeUnregistered) {
+    // pointer not known to the HIP subsystem
+    return std::nullopt;
+  }
+#endif
+  return hipPointerAttributes;
+}
+
+// Helper method to abstract away the fact that retrieving a pointer's memory
+// type differs depending on the version of HIP.
+inline static unsigned getMemoryType(hipPointerAttribute_t hipPointerAttrs) {
+#if HIP_VERSION >= 50600000
+  return hipPointerAttrs.type;
+#else
+  return hipPointerAttrs.memoryType;
+#endif
+}

source/adapters/hip/device.cpp

@@ -782,8 +782,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
         UR_MEMORY_ORDER_CAPABILITY_FLAG_RELEASE;
     return ReturnValue(Capabilities);
   }
-  case UR_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES:
-  case UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES: {
+  case UR_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES: {
     // SYCL2020 4.6.4.2 minimum mandated capabilities for
     // atomic_fence/memory_scope_capabilities.
     // Because scopes are hierarchical, wider scopes support all narrower
@@ -795,6 +794,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
         UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_GROUP;
     return ReturnValue(Capabilities);
   }
+  case UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES: {
+    constexpr ur_memory_scope_capability_flags_t Capabilities =
+        UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_ITEM |
+        UR_MEMORY_SCOPE_CAPABILITY_FLAG_SUB_GROUP |
+        UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_GROUP |
+        UR_MEMORY_SCOPE_CAPABILITY_FLAG_DEVICE |
+        UR_MEMORY_SCOPE_CAPABILITY_FLAG_SYSTEM;
+    return ReturnValue(Capabilities);
+  }
   case UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: {
     // SYCL2020 4.6.4.2 minimum mandated capabilities for
     // atomic_fence_order_capabilities.
@@ -803,6 +811,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
         UR_MEMORY_ORDER_CAPABILITY_FLAG_ACQUIRE |
         UR_MEMORY_ORDER_CAPABILITY_FLAG_RELEASE |
         UR_MEMORY_ORDER_CAPABILITY_FLAG_ACQ_REL;
+#ifdef __HIP_PLATFORM_NVIDIA__
+    // Nvidia introduced fence.sc for seq_cst only since SM 7.0.
+    int Major = 0;
+    UR_CHECK_ERROR(hipDeviceGetAttribute(
+        &Major, hipDeviceAttributeComputeCapabilityMajor, hDevice->get()));
+    if (Major >= 7)
+      Capabilities |= UR_MEMORY_ORDER_CAPABILITY_FLAG_SEQ_CST;
+#else
+    Capabilities |= UR_MEMORY_ORDER_CAPABILITY_FLAG_SEQ_CST;
+#endif
     return ReturnValue(Capabilities);
   }
   case UR_DEVICE_INFO_DEVICE_ID: {

source/adapters/hip/enqueue.cpp

@@ -789,21 +789,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferFill(
     ur_event_handle_t *phEvent) {
   UR_ASSERT(size + offset <= std::get<BufferMem>(hBuffer->Mem).getSize(),
             UR_RESULT_ERROR_INVALID_SIZE);
-  auto ArgsAreMultiplesOfPatternSize =
-      (offset % patternSize == 0) || (size % patternSize == 0);
-
-  auto PatternIsValid = (pPattern != nullptr);
-
-  auto PatternSizeIsValid =
-      ((patternSize & (patternSize - 1)) == 0) && // is power of two
-      (patternSize > 0) && (patternSize <= 128);  // falls within valid range
-
-  UR_ASSERT(ArgsAreMultiplesOfPatternSize && PatternIsValid &&
-                PatternSizeIsValid,
-            UR_RESULT_ERROR_INVALID_VALUE);
-  std::ignore = ArgsAreMultiplesOfPatternSize;
-  std::ignore = PatternIsValid;
-  std::ignore = PatternSizeIsValid;
 
   std::unique_ptr<ur_event_handle_t_> RetImplEvent{nullptr};
   hBuffer->setLastQueueWritingToMemObj(hQueue);

source/adapters/hip/program.cpp

@@ -283,7 +283,11 @@ urProgramCreateWithIL(ur_context_handle_t, const void *, size_t,
 UR_APIEXPORT ur_result_t UR_APICALL
 urProgramCompile(ur_context_handle_t hContext, ur_program_handle_t hProgram,
                  const char *pOptions) {
-  return urProgramBuild(hContext, hProgram, pOptions);
+  UR_CHECK_ERROR(urProgramBuild(hContext, hProgram, pOptions));
+  // urProgramBuild sets the BinaryType to UR_PROGRAM_BINARY_TYPE_EXECUTABLE, so
+  // set it to the correct value for urProgramCompile post-hoc.
+  hProgram->BinaryType = UR_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
+  return UR_RESULT_SUCCESS;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urProgramCompileExp(ur_program_handle_t,
@@ -312,6 +316,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramBuild(ur_context_handle_t,
     ScopedContext Active(hProgram->getDevice());
 
     hProgram->buildProgram(pOptions);
+    hProgram->BinaryType = UR_PROGRAM_BINARY_TYPE_EXECUTABLE;
 
   } catch (ur_result_t Err) {
     Result = Err;
@@ -355,13 +360,14 @@ urProgramGetBuildInfo(ur_program_handle_t hProgram, ur_device_handle_t,
   UrReturnHelper ReturnValue(propSize, pPropValue, pPropSizeRet);
 
   switch (propName) {
-  case UR_PROGRAM_BUILD_INFO_STATUS: {
+  case UR_PROGRAM_BUILD_INFO_STATUS:
     return ReturnValue(hProgram->BuildStatus);
-  }
   case UR_PROGRAM_BUILD_INFO_OPTIONS:
     return ReturnValue(hProgram->BuildOptions.c_str());
   case UR_PROGRAM_BUILD_INFO_LOG:
     return ReturnValue(hProgram->InfoLog, hProgram->MAX_LOG_SIZE);
+  case UR_PROGRAM_BUILD_INFO_BINARY_TYPE:
+    return ReturnValue(hProgram->BinaryType);
   default:
     break;
   }
@@ -494,6 +500,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithBinary(
   UR_ASSERT(Result == UR_RESULT_SUCCESS, Result);
 
   *phProgram = RetProgram.release();
+  (*phProgram)->BinaryType = UR_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
 
   return Result;
 }

source/adapters/hip/program.hpp

@@ -27,6 +27,12 @@ struct ur_program_handle_t_ {
   ur_device_handle_t Device;
   std::string ExecutableCache;
 
+  // The ur_program_binary_type_t property is defined individually for every
+  // device in a program. However, since the HIP adapter only has 1 device per
+  // program, there is no need to keep track of its value for each
+  // device.
+  ur_program_binary_type_t BinaryType = UR_PROGRAM_BINARY_TYPE_NONE;
+
   // Metadata
   bool IsRelocatable = false;
 

source/adapters/hip/queue.cpp

@@ -191,10 +191,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueGetInfo(ur_queue_handle_t hQueue,
     });
     return ReturnValue(IsReady);
   }
+  case UR_QUEUE_INFO_DEVICE_DEFAULT:
+  case UR_QUEUE_INFO_SIZE:
+    return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
   default:
-    break;
+    return UR_RESULT_ERROR_INVALID_ENUMERATION;
   }
-  return {};
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urQueueRetain(ur_queue_handle_t hQueue) {

source/adapters/hip/usm.cpp

@@ -152,41 +152,21 @@ UR_APIEXPORT ur_result_t UR_APICALL
 urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem,
                      ur_usm_alloc_info_t propName, size_t propValueSize,
                      void *pPropValue, size_t *pPropValueSizeRet) {
-  ur_result_t Result = UR_RESULT_SUCCESS;
-  hipPointerAttribute_t hipPointerAttributeType;
-
   UrReturnHelper ReturnValue(propValueSize, pPropValue, pPropValueSizeRet);
 
   try {
     switch (propName) {
     case UR_USM_ALLOC_INFO_TYPE: {
-      // do not throw if hipPointerGetAttribute returns hipErrorInvalidValue
-      hipError_t Ret = hipPointerGetAttributes(&hipPointerAttributeType, pMem);
-      if (Ret == hipErrorInvalidValue) {
-        // pointer not known to the HIP subsystem
-        return ReturnValue(UR_USM_TYPE_UNKNOWN);
-      }
-      // Direct usage of the function, instead of UR_CHECK_ERROR, so we can get
-      // the line offset.
-      checkErrorUR(Ret, __func__, __LINE__ - 5, __FILE__);
-      // ROCm 6.0.0 introduces hipMemoryTypeUnregistered in the hipMemoryType
-      // enum to mark unregistered allocations (i.e., via system allocators).
-#if HIP_VERSION_MAJOR >= 6
-      if (hipPointerAttributeType.type == hipMemoryTypeUnregistered) {
+      auto MaybePointerAttrs = getPointerAttributes(pMem);
+      if (!MaybePointerAttrs.has_value()) {
         // pointer not known to the HIP subsystem
         return ReturnValue(UR_USM_TYPE_UNKNOWN);
       }
-#endif
-      unsigned int Value;
-#if HIP_VERSION >= 50600000
-      Value = hipPointerAttributeType.type;
-#else
-      Value = hipPointerAttributeType.memoryType;
-#endif
+      auto Value = getMemoryType(*MaybePointerAttrs);
       UR_ASSERT(Value == hipMemoryTypeDevice || Value == hipMemoryTypeHost ||
                     Value == hipMemoryTypeManaged,
                 UR_RESULT_ERROR_INVALID_MEM_OBJECT);
-      if (hipPointerAttributeType.isManaged || Value == hipMemoryTypeManaged) {
+      if (MaybePointerAttrs->isManaged || Value == hipMemoryTypeManaged) {
         // pointer to managed memory
         return ReturnValue(UR_USM_TYPE_SHARED);
       }
@@ -202,15 +182,18 @@ urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem,
       ur::unreachable();
     }
     case UR_USM_ALLOC_INFO_DEVICE: {
-      // get device index associated with this pointer
-      UR_CHECK_ERROR(hipPointerGetAttributes(&hipPointerAttributeType, pMem));
+      auto MaybePointerAttrs = getPointerAttributes(pMem);
+      if (!MaybePointerAttrs.has_value()) {
+        // pointer not known to the HIP subsystem
+        return ReturnValue(UR_USM_TYPE_UNKNOWN);
+      }
 
-      int DeviceIdx = hipPointerAttributeType.device;
+      int DeviceIdx = MaybePointerAttrs->device;
 
       // hip backend has only one platform containing all devices
       ur_platform_handle_t platform;
       ur_adapter_handle_t AdapterHandle = &adapter;
-      Result = urPlatformGet(&AdapterHandle, 1, 1, &platform, nullptr);
+      UR_CHECK_ERROR(urPlatformGet(&AdapterHandle, 1, 1, &platform, nullptr));
 
       // get the device from the platform
       ur_device_handle_t Device = platform->Devices[DeviceIdx].get();
@@ -227,20 +210,32 @@ urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem,
       }
       return ReturnValue(Pool);
     }
+    case UR_USM_ALLOC_INFO_BASE_PTR:
+      // HIP gives us the ability to query the base pointer for a device
+      // pointer, so check whether we've got one of those.
+      if (auto MaybePointerAttrs = getPointerAttributes(pMem)) {
+        if (getMemoryType(*MaybePointerAttrs) == hipMemoryTypeDevice) {
+          void *Base = nullptr;
+          UR_CHECK_ERROR(hipPointerGetAttribute(
+              &Base, HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR,
+              (hipDeviceptr_t)pMem));
+          return ReturnValue(Base);
+        }
+      }
+      // If not, we can't be sure.
+      return UR_RESULT_ERROR_INVALID_VALUE;
     case UR_USM_ALLOC_INFO_SIZE: {
       size_t RangeSize = 0;
       UR_CHECK_ERROR(hipMemPtrGetInfo(const_cast<void *>(pMem), &RangeSize));
       return ReturnValue(RangeSize);
     }
-    case UR_USM_ALLOC_INFO_BASE_PTR:
-      return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION;
     default:
       return UR_RESULT_ERROR_INVALID_ENUMERATION;
     }
   } catch (ur_result_t Error) {
-    Result = Error;
+    return Error;
   }
-  return Result;
+  return UR_RESULT_SUCCESS;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urUSMImportExp(ur_context_handle_t Context,

test/conformance/device_code/cpy_and_mult.cpp

@@ -7,20 +7,20 @@
 
 int main() {
     size_t array_size = 16;
-    cl::sycl::queue sycl_queue;
+    sycl::queue sycl_queue;
     std::vector<uint32_t> src(array_size, 1);
     std::vector<uint32_t> dst(array_size, 1);
     auto src_buff =
-        cl::sycl::buffer<uint32_t>(src.data(), cl::sycl::range<1>(array_size));
+        sycl::buffer<uint32_t>(src.data(), sycl::range<1>(array_size));
     auto dst_buff =
-        cl::sycl::buffer<uint32_t>(dst.data(), cl::sycl::range<1>(array_size));
+        sycl::buffer<uint32_t>(dst.data(), sycl::range<1>(array_size));
 
-    sycl_queue.submit([&](cl::sycl::handler &cgh) {
-        auto src_acc = src_buff.get_access<cl::sycl::access::mode::read>(cgh);
-        auto dst_acc = dst_buff.get_access<cl::sycl::access::mode::write>(cgh);
+    sycl_queue.submit([&](sycl::handler &cgh) {
+        auto src_acc = src_buff.get_access<sycl::access::mode::read>(cgh);
+        auto dst_acc = dst_buff.get_access<sycl::access::mode::write>(cgh);
         cgh.parallel_for<class cpy_and_mult>(
-            cl::sycl::range<1>{array_size},
-            [src_acc, dst_acc](cl::sycl::item<1> itemId) {
+            sycl::range<1>{array_size},
+            [src_acc, dst_acc](sycl::item<1> itemId) {
                 auto id = itemId.get_id(0);
                 dst_acc[id] = src_acc[id] * 2;
             });

test/conformance/device_code/cpy_and_mult_usm.cpp

@@ -7,13 +7,12 @@
 
 int main() {
     size_t array_size = 16;
-    cl::sycl::queue sycl_queue;
-    uint32_t *src = cl::sycl::malloc_device<uint32_t>(array_size, sycl_queue);
-    uint32_t *dst = cl::sycl::malloc_device<uint32_t>(array_size, sycl_queue);
-    sycl_queue.submit([&](cl::sycl::handler &cgh) {
+    sycl::queue sycl_queue;
+    uint32_t *src = sycl::malloc_device<uint32_t>(array_size, sycl_queue);
+    uint32_t *dst = sycl::malloc_device<uint32_t>(array_size, sycl_queue);
+    sycl_queue.submit([&](sycl::handler &cgh) {
         cgh.parallel_for<class cpy_and_mult_usm>(
-            cl::sycl::range<1>{array_size},
-            [src, dst](cl::sycl::item<1> itemId) {
+            sycl::range<1>{array_size}, [src, dst](sycl::item<1> itemId) {
                 auto id = itemId.get_id(0);
                 dst[id] = src[id] * 2;
             });

test/conformance/enqueue/enqueue_adapter_hip.match

@@ -5,8 +5,6 @@ urEnqueueKernelLaunchKernelSubGroupTest.Success/AMD_HIP_BACKEND___{{.*}}_
 urEnqueueKernelLaunchUSMLinkedList.Success/AMD_HIP_BACKEND___{{.*}}___UsePoolEnabled
 {{OPT}}urEnqueueMemBufferCopyRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___copy_row_2D
 {{OPT}}urEnqueueMemBufferCopyRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___copy_3d_2d
-urEnqueueMemBufferFillTest.Success/AMD_HIP_BACKEND___{{.*}}___size__256__patternSize__256
-urEnqueueMemBufferFillTest.Success/AMD_HIP_BACKEND___{{.*}}___size__1024__patternSize__256
 {{OPT}}urEnqueueMemBufferWriteRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___write_row_2D
 {{OPT}}urEnqueueMemBufferWriteRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___write_3d_2d
 urEnqueueUSMAdviseWithParamTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_USM_ADVICE_FLAG_DEFAULT