[SYCL] Adds info query for atomic_memory_order_capabilities on device and context

sycl/include/CL/sycl/detail/pi.h

@@ -292,7 +292,8 @@ typedef enum {
   PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE = 0x10025,
   PI_DEVICE_INFO_MAX_MEM_BANDWIDTH = 0x10026,
   PI_DEVICE_INFO_IMAGE_SRGB = 0x10027,
-  PI_DEVICE_INFO_ATOMIC_64 = 0x10110
+  PI_DEVICE_INFO_ATOMIC_64 = 0x10110,
+  PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES = 0x10111
 } _pi_device_info;
 
 typedef enum {
@@ -312,6 +313,8 @@ typedef enum {
   PI_CONTEXT_INFO_NUM_DEVICES = CL_CONTEXT_NUM_DEVICES,
   PI_CONTEXT_INFO_PROPERTIES = CL_CONTEXT_PROPERTIES,
   PI_CONTEXT_INFO_REFERENCE_COUNT = CL_CONTEXT_REFERENCE_COUNT,
+  // Atomics capabilities extensions
+  PI_CONTEXT_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES = 0x10010
 } _pi_context_info;
 
 typedef enum {
@@ -509,6 +512,13 @@ constexpr pi_sampler_properties PI_SAMPLER_PROPERTIES_ADDRESSING_MODE =
 constexpr pi_sampler_properties PI_SAMPLER_PROPERTIES_FILTER_MODE =
     CL_SAMPLER_FILTER_MODE;
 
+using pi_memory_order_capabilities = pi_bitfield;
+constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_RELAXED = 0x01;
+constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_ACQUIRE = 0x02;
+constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_RELEASE = 0x04;
+constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_ACQ_REL = 0x08;
+constexpr pi_memory_order_capabilities PI_MEMORY_ORDER_SEQ_CST = 0x10;
+
 typedef enum {
   PI_PROFILING_INFO_COMMAND_QUEUED = CL_PROFILING_COMMAND_QUEUED,
   PI_PROFILING_INFO_COMMAND_SUBMIT = CL_PROFILING_COMMAND_SUBMIT,

sycl/include/CL/sycl/detail/spirv.hpp

@@ -224,11 +224,10 @@ EnableIfGenericBroadcast<T> GroupBroadcast(T x, id<Dimensions> local_id) {
 // Single happens-before means semantics should always apply to all spaces
 // Although consume is unsupported, forwarding to acquire is valid
 template <typename T>
-static inline constexpr typename std::enable_if<
-    std::is_same<T, sycl::ext::oneapi::memory_order>::value ||
-        std::is_same<T, sycl::memory_order>::value,
-    __spv::MemorySemanticsMask::Flag>::type
-getMemorySemanticsMask(T Order) {
+static inline constexpr
+    typename std::enable_if<std::is_same<T, sycl::memory_order>::value,
+                            __spv::MemorySemanticsMask::Flag>::type
+    getMemorySemanticsMask(T Order) {
   __spv::MemorySemanticsMask::Flag SpvOrder = __spv::MemorySemanticsMask::None;
   switch (Order) {
   case T::relaxed:
@@ -254,28 +253,25 @@ getMemorySemanticsMask(T Order) {
       __spv::MemorySemanticsMask::CrossWorkgroupMemory);
 }
 
-static inline constexpr __spv::Scope::Flag
-getScope(ext::oneapi::memory_scope Scope) {
+static inline constexpr __spv::Scope::Flag getScope(memory_scope Scope) {
   switch (Scope) {
-  case ext::oneapi::memory_scope::work_item:
+  case memory_scope::work_item:
     return __spv::Scope::Invocation;
-  case ext::oneapi::memory_scope::sub_group:
+  case memory_scope::sub_group:
     return __spv::Scope::Subgroup;
-  case ext::oneapi::memory_scope::work_group:
+  case memory_scope::work_group:
     return __spv::Scope::Workgroup;
-  case ext::oneapi::memory_scope::device:
+  case memory_scope::device:
     return __spv::Scope::Device;
-  case ext::oneapi::memory_scope::system:
+  case memory_scope::system:
     return __spv::Scope::CrossDevice;
   }
 }
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
-                      ext::oneapi::memory_scope Scope,
-                      ext::oneapi::memory_order Success,
-                      ext::oneapi::memory_order Failure, T Desired,
+AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+                      memory_order Success, memory_order Failure, T Desired,
                       T Expected) {
   auto SPIRVSuccess = getMemorySemanticsMask(Success);
   auto SPIRVFailure = getMemorySemanticsMask(Failure);
@@ -287,10 +283,8 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
-AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
-                      ext::oneapi::memory_scope Scope,
-                      ext::oneapi::memory_order Success,
-                      ext::oneapi::memory_order Failure, T Desired,
+AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+                      memory_order Success, memory_order Failure, T Desired,
                       T Expected) {
   using I = detail::make_unsinged_integer_t<T>;
   auto SPIRVSuccess = getMemorySemanticsMask(Success);
@@ -308,8 +302,8 @@ AtomicCompareExchange(multi_ptr<T, AddressSpace> MPtr,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-           ext::oneapi::memory_order Order) {
+AtomicLoad(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+           memory_order Order) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -318,8 +312,8 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
-AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-           ext::oneapi::memory_order Order) {
+AtomicLoad(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+           memory_order Order) {
   using I = detail::make_unsinged_integer_t<T>;
   auto *PtrInt =
       reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -332,8 +326,8 @@ AtomicLoad(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value>
-AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-            ext::oneapi::memory_order Order, T Value) {
+AtomicStore(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+            memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -342,8 +336,8 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value>
-AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-            ext::oneapi::memory_order Order, T Value) {
+AtomicStore(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+            memory_order Order, T Value) {
   using I = detail::make_unsinged_integer_t<T>;
   auto *PtrInt =
       reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -356,8 +350,8 @@ AtomicStore(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-               ext::oneapi::memory_order Order, T Value) {
+AtomicExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+               memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -366,8 +360,8 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
-AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-               ext::oneapi::memory_order Order, T Value) {
+AtomicExchange(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+               memory_order Order, T Value) {
   using I = detail::make_unsinged_integer_t<T>;
   auto *PtrInt =
       reinterpret_cast<typename multi_ptr<I, AddressSpace>::pointer_t>(
@@ -382,8 +376,8 @@ AtomicExchange(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-           ext::oneapi::memory_order Order, T Value) {
+AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+           memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -392,8 +386,8 @@ AtomicIAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicISub(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-           ext::oneapi::memory_order Order, T Value) {
+AtomicISub(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+           memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -402,8 +396,8 @@ AtomicISub(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
-AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-           ext::oneapi::memory_order Order, T Value) {
+AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+           memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -412,8 +406,8 @@ AtomicFAdd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicAnd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-          ext::oneapi::memory_order Order, T Value) {
+AtomicAnd(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -422,8 +416,8 @@ AtomicAnd(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicOr(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-         ext::oneapi::memory_order Order, T Value) {
+AtomicOr(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+         memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -432,8 +426,8 @@ AtomicOr(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicXor(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-          ext::oneapi::memory_order Order, T Value) {
+AtomicXor(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -442,8 +436,8 @@ AtomicXor(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-          ext::oneapi::memory_order Order, T Value) {
+AtomicMin(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -452,8 +446,8 @@ AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
-AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-          ext::oneapi::memory_order Order, T Value) {
+AtomicMin(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -462,8 +456,8 @@ AtomicMin(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_integral<T>::value, T>
-AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-          ext::oneapi::memory_order Order, T Value) {
+AtomicMax(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -472,8 +466,8 @@ AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
 
 template <typename T, access::address_space AddressSpace>
 inline typename detail::enable_if_t<std::is_floating_point<T>::value, T>
-AtomicMax(multi_ptr<T, AddressSpace> MPtr, ext::oneapi::memory_scope Scope,
-          ext::oneapi::memory_order Order, T Value) {
+AtomicMax(multi_ptr<T, AddressSpace> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
   auto *Ptr = MPtr.get();
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);

sycl/include/CL/sycl/info/context_traits.def

@@ -1,3 +1,4 @@
 __SYCL_PARAM_TRAITS_SPEC(context, reference_count, cl_uint)
 __SYCL_PARAM_TRAITS_SPEC(context, platform, cl::sycl::platform)
 __SYCL_PARAM_TRAITS_SPEC(context, devices, std::vector<cl::sycl::device>)
+__SYCL_PARAM_TRAITS_SPEC(context, atomic_memory_order_capabilities, std::vector<cl::sycl::memory_order>)

sycl/include/CL/sycl/info/device_traits.def

@@ -23,6 +23,8 @@ __SYCL_PARAM_TRAITS_SPEC(device, address_bits, pi_uint32)
 __SYCL_PARAM_TRAITS_SPEC(device, max_mem_alloc_size, pi_uint64)
 __SYCL_PARAM_TRAITS_SPEC(device, image_support, bool)
 __SYCL_PARAM_TRAITS_SPEC(device, atomic64, bool)
+__SYCL_PARAM_TRAITS_SPEC(device, atomic_memory_order_capabilities,
+                         std::vector<cl::sycl::memory_order>)
 __SYCL_PARAM_TRAITS_SPEC(device, max_read_image_args, pi_uint32)
 __SYCL_PARAM_TRAITS_SPEC(device, max_write_image_args, pi_uint32)
 __SYCL_PARAM_TRAITS_SPEC(device, image2d_max_width, size_t)

sycl/include/CL/sycl/info/info_desc.hpp

@@ -37,6 +37,8 @@ enum class context : cl_context_info {
   reference_count = CL_CONTEXT_REFERENCE_COUNT,
   platform = CL_CONTEXT_PLATFORM,
   devices = CL_CONTEXT_DEVICES,
+  atomic_memory_order_capabilities =
+      PI_CONTEXT_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES,
 };
 
 // A.3 Device information descriptors
@@ -149,7 +151,9 @@ enum class device : cl_device_info {
   ext_intel_mem_channel = PI_MEM_PROPERTIES_CHANNEL,
   ext_oneapi_srgb = PI_DEVICE_INFO_IMAGE_SRGB,
   ext_intel_device_info_uuid = PI_DEVICE_INFO_UUID,
-  atomic64 = PI_DEVICE_INFO_ATOMIC_64
+  atomic64 = PI_DEVICE_INFO_ATOMIC_64,
+  atomic_memory_order_capabilities =
+      PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES,
 };
 
 enum class device_type : pi_uint64 {

sycl/include/CL/sycl/memory_enums.hpp

@@ -8,11 +8,8 @@
 
 #pragma once
 
-#include <sycl/ext/oneapi/atomic_enums.hpp>
-
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-using ext::oneapi::memory_scope;
 
 enum class memory_order : int {
   relaxed = 0,
@@ -24,6 +21,14 @@ enum class memory_order : int {
   seq_cst = 5
 };
 
+enum class memory_scope : int {
+  work_item = 0,
+  sub_group = 1,
+  work_group = 2,
+  device = 3,
+  system = 4
+};
+
 #if __cplusplus >= 201703L
 inline constexpr auto memory_scope_work_item = memory_scope::work_item;
 inline constexpr auto memory_scope_sub_group = memory_scope::sub_group;
@@ -38,9 +43,25 @@ inline constexpr auto memory_order_acq_rel = memory_order::acq_rel;
 inline constexpr auto memory_order_seq_cst = memory_order::seq_cst;
 #endif
 
-#ifndef __SYCL_DEVICE_ONLY__
 namespace detail {
 
+inline std::vector<memory_order>
+readMemoryOrderBitfield(pi_memory_order_capabilities bits) {
+  std::vector<memory_order> result;
+  if (bits & PI_MEMORY_ORDER_RELAXED)
+    result.push_back(memory_order::relaxed);
+  if (bits & PI_MEMORY_ORDER_ACQUIRE)
+    result.push_back(memory_order::acquire);
+  if (bits & PI_MEMORY_ORDER_RELEASE)
+    result.push_back(memory_order::release);
+  if (bits & PI_MEMORY_ORDER_ACQ_REL)
+    result.push_back(memory_order::acq_rel);
+  if (bits & PI_MEMORY_ORDER_SEQ_CST)
+    result.push_back(memory_order::seq_cst);
+  return result;
+}
+
+#ifndef __SYCL_DEVICE_ONLY__
 static constexpr std::memory_order getStdMemoryOrder(sycl::memory_order order) {
   switch (order) {
   case memory_order::relaxed:
@@ -57,8 +78,8 @@ static constexpr std::memory_order getStdMemoryOrder(sycl::memory_order order) {
     return std::memory_order_seq_cst;
   }
 }
+#endif // __SYCL_DEVICE_ONLY__
 
 } // namespace detail
-#endif // __SYCL_DEVICE_ONLY__
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)