[SYCL] Adds support for atomic fence capabilities device queries

opencl/CMakeLists.txt

@@ -44,7 +44,7 @@ FetchContent_GetProperties(ocl-headers)
 set(OpenCL_INCLUDE_DIR
   ${ocl-headers_SOURCE_DIR} CACHE PATH "Path to OpenCL Headers")
 
-target_compile_definitions(Headers INTERFACE -DCL_TARGET_OPENCL_VERSION=220)
+target_compile_definitions(Headers INTERFACE -DCL_TARGET_OPENCL_VERSION=300)
 add_library(OpenCL-Headers ALIAS Headers)
 
 # OpenCL Library (ICD Loader)

sycl/include/sycl/detail/cl.h

@@ -9,9 +9,9 @@
 #pragma once
 
 // Suppress a compiler message about undefined CL_TARGET_OPENCL_VERSION
-// and define all symbols up to OpenCL 2.2
+// and define all symbols up to OpenCL 3.0
 #ifndef CL_TARGET_OPENCL_VERSION
-#define CL_TARGET_OPENCL_VERSION 220
+#define CL_TARGET_OPENCL_VERSION 300
 #endif
 
 #include <CL/cl.h>

sycl/include/sycl/detail/pi.h

@@ -77,9 +77,11 @@
 // 12.22 Add piGetDeviceAndHostTimer to query device wall-clock timestamp
 // 12.23 Added new piextEnqueueDeviceGlobalVariableWrite and
 // piextEnqueueDeviceGlobalVariableRead functions.
+// 12.24 Added PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES and
+// PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES for piDeviceGetInfo.
 
 #define _PI_H_VERSION_MAJOR 12
-#define _PI_H_VERSION_MINOR 23
+#define _PI_H_VERSION_MINOR 24
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -313,6 +315,8 @@ typedef enum {
   PI_DEVICE_INFO_ATOMIC_64 = 0x10110,
   PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES = 0x10111,
   PI_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES = 0x11000,
+  PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES = 0x10114,
+  PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES = 0x10115,
   PI_DEVICE_INFO_GPU_HW_THREADS_PER_EU = 0x10112,
   PI_DEVICE_INFO_BACKEND_VERSION = 0x10113,
   // Return whether bfloat16 math functions are supported by device

sycl/include/sycl/info/device_traits.def

@@ -112,9 +112,15 @@ __SYCL_PARAM_TRAITS_SPEC(device, host_unified_memory, bool,
 __SYCL_PARAM_TRAITS_SPEC(device, atomic_memory_order_capabilities,
                          std::vector<sycl::memory_order>,
                          PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES)
+__SYCL_PARAM_TRAITS_SPEC(device, atomic_fence_order_capabilities,
+                         std::vector<sycl::memory_order>,
+                         PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES)
 __SYCL_PARAM_TRAITS_SPEC(device, atomic_memory_scope_capabilities,
                          std::vector<sycl::memory_scope>,
                          PI_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES)
+__SYCL_PARAM_TRAITS_SPEC(device, atomic_fence_scope_capabilities,
+                         std::vector<sycl::memory_scope>,
+                         PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES)
 __SYCL_PARAM_TRAITS_SPEC(device, profiling_timer_resolution, size_t,
                          PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION)
 __SYCL_PARAM_TRAITS_SPEC(device, is_endian_little, bool,

sycl/include/sycl/info/info_desc.hpp

@@ -95,6 +95,9 @@ namespace device {
 // atomic_fence_order_capabilities, atomic_fence_scope_capabilities, aspects,
 // il_version.
 
+struct atomic_fence_order_capabilities;
+struct atomic_fence_scope_capabilities;
+
 #define __SYCL_PARAM_TRAITS_DEPRECATED(Desc, Message)                          \
   struct __SYCL2020_DEPRECATED(Message) Desc;
 #include <sycl/info/device_traits_deprecated.def>

sycl/plugins/cuda/pi_cuda.cpp

@@ -1315,6 +1315,10 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
     return getInfo(param_value_size, param_value, param_value_size_ret,
                    capabilities);
   }
+  case PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES:
+  case PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES:
+    // There is no way to query this in the backend
+    return PI_ERROR_INVALID_ARG_VALUE;
   case PI_EXT_ONEAPI_DEVICE_INFO_BFLOAT16_MATH_FUNCTIONS: {
     int major = 0;
     sycl::detail::pi::assertion(

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -805,6 +805,8 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
     CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_IMAGE_SRGB)
     CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_ATOMIC_64)
     CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES)
+    CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES)
+    CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES)
     CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS)
     CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D)
     CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D)

sycl/plugins/hip/pi_hip.cpp

@@ -529,9 +529,9 @@ hipStream_t _pi_queue::get_next_transfer_stream() {
 _pi_event::_pi_event(pi_command_type type, pi_context context, pi_queue queue,
                      hipStream_t stream, pi_uint32 stream_token)
     : commandType_{type}, refCount_{1}, hasBeenWaitedOn_{false},
-      isRecorded_{false}, isStarted_{false},
-      streamToken_{stream_token}, evEnd_{nullptr}, evStart_{nullptr},
-      evQueued_{nullptr}, queue_{queue}, stream_{stream}, context_{context} {
+      isRecorded_{false}, isStarted_{false}, streamToken_{stream_token},
+      evEnd_{nullptr}, evStart_{nullptr}, evQueued_{nullptr}, queue_{queue},
+      stream_{stream}, context_{context} {
 
   assert(type != PI_COMMAND_TYPE_USER);
 
@@ -685,8 +685,8 @@ pi_result enqueueEventWait(pi_queue queue, pi_event event) {
 }
 
 _pi_program::_pi_program(pi_context ctxt)
-    : module_{nullptr}, binary_{},
-      binarySizeInBytes_{0}, refCount_{1}, context_{ctxt} {
+    : module_{nullptr}, binary_{}, binarySizeInBytes_{0}, refCount_{1},
+      context_{ctxt} {
   hip_piContextRetain(context_);
 }
 
@@ -1865,6 +1865,9 @@ pi_result hip_piDeviceGetInfo(pi_device device, pi_device_info param_name,
   }
   // TODO: Investigate if this information is available on HIP.
   case PI_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES:
+  // There is no way to query this in the backend
+  case PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES:
+  case PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES:
   case PI_DEVICE_INFO_DEVICE_ID:
   case PI_DEVICE_INFO_PCI_ADDRESS:
   case PI_DEVICE_INFO_GPU_EU_COUNT:
@@ -5324,7 +5327,7 @@ pi_result hip_piextEnqueueDeviceGlobalVariableRead(
 // Windows: dynamically loaded plugins might have been unloaded already
 // when this is called. Sycl RT holds onto the PI plugin so it can be
 // called safely. But this is not transitive. If the PI plugin in turn
-// dynamically loaded a different DLL, that may have been unloaded. 
+// dynamically loaded a different DLL, that may have been unloaded.
 // TODO: add a global variable lifetime management code here (see
 // pi_level_zero.cpp for reference) Currently this is just a NOOP.
 pi_result hip_piTearDown(void *PluginParameter) {

sycl/plugins/opencl/pi_opencl.cpp

@@ -285,6 +285,106 @@ pi_result piDeviceGetInfo(pi_device device, pi_device_info paramName,
   case PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES:
   case PI_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES:
     return PI_ERROR_INVALID_VALUE;
+  case PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: {
+    // Initialize result to minimum mandated capabilities according to
+    // SYCL2020 4.6.3.2
+    pi_memory_order_capabilities result =
+        PI_MEMORY_ORDER_RELAXED | PI_MEMORY_ORDER_ACQUIRE |
+        PI_MEMORY_ORDER_RELEASE | PI_MEMORY_ORDER_ACQ_REL;
+
+    OCLV::OpenCLVersion devVer;
+
+    cl_device_id deviceID = cast<cl_device_id>(device);
+    cl_int ret_err = getDeviceVersion(deviceID, devVer);
+    if (ret_err != CL_SUCCESS)
+      return static_cast<pi_result>(ret_err);
+
+    cl_device_atomic_capabilities devCapabilities = 0;
+    if (devVer >= OCLV::V3_0) {
+      ret_err = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_FENCE_CAPABILITIES,
+                                sizeof(cl_device_atomic_capabilities),
+                                &devCapabilities, nullptr);
+      if (ret_err != CL_SUCCESS)
+        return static_cast<pi_result>(ret_err);
+      assert((devCapabilities & CL_DEVICE_ATOMIC_ORDER_RELAXED) &&
+             "Violates minimum mandated guarantee");
+      assert((devCapabilities & CL_DEVICE_ATOMIC_ORDER_ACQ_REL) &&
+             "Violates minimum mandated guarantee");
+
+      if (devCapabilities & CL_DEVICE_ATOMIC_ORDER_SEQ_CST) {
+        result |= PI_MEMORY_ORDER_SEQ_CST;
+      }
+
+    } else {
+      // This info is only available in OpenCL version >= 3.0
+      // Just return minimum mandated capabilities for older versions.
+      // OpenCL 1.x minimum mandated capabilities are RELAXED | ACQ_REL, we
+      // already initialized using these.
+      if (devVer >= OCLV::V2_0) {
+        // OpenCL 2.x minimum mandated capabilities are RELAXED | ACQ_REL |
+        // SEQ_CST
+        result |= PI_MEMORY_ORDER_SEQ_CST;
+      }
+    }
+    std::memcpy(paramValue, &result, sizeof(result));
+    return PI_SUCCESS;
+  }
+  case PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES: {
+    // Initialize result to minimum mandated capabilities according to
+    // SYCL2020 4.6.3.2.
+    pi_memory_scope_capabilities result = PI_MEMORY_SCOPE_WORK_ITEM |
+                                          PI_MEMORY_SCOPE_SUB_GROUP |
+                                          PI_MEMORY_SCOPE_WORK_GROUP;
+
+    OCLV::OpenCLVersion devVer;
+
+    cl_device_id deviceID = cast<cl_device_id>(device);
+    cl_int ret_err = getDeviceVersion(deviceID, devVer);
+    if (ret_err != CL_SUCCESS)
+      return static_cast<pi_result>(ret_err);
+
+    cl_device_atomic_capabilities devCapabilities = 0;
+    if (devVer >= OCLV::V3_0) {
+      ret_err = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_FENCE_CAPABILITIES,
+                                sizeof(cl_device_atomic_capabilities),
+                                &devCapabilities, nullptr);
+      if (ret_err != CL_SUCCESS)
+        return static_cast<pi_result>(ret_err);
+      assert((devCapabilities & CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP) &&
+             "Violates minimum mandated guarantee");
+
+      // Because scopes are hierarchical, wider scopes support all narrower
+      // scopes. SUB_GROUP and WORK_ITEM was already included in the
+      // initialization, since WORK_GROUP is mandated minimum capality.
+      if (devCapabilities & CL_DEVICE_ATOMIC_SCOPE_DEVICE) {
+        result |= PI_MEMORY_SCOPE_DEVICE;
+      }
+
+      if (devCapabilities & CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES) {
+        result |= PI_MEMORY_SCOPE_SYSTEM;
+      }
+
+    } else {
+      // This info is only available in OpenCL version >= 3.0
+      // Just return minimum mandated capabilities for older versions.
+      // OpenCL 1.x minimum mandated capabilities are WORK_GROUP, we
+      // already initialized using it.
+      if (devVer >= OCLV::V2_0) {
+        // OpenCL 2.x minimum mandated capabilities are WORK_GROUP | DEVICE |
+        // ALL_DEVICES
+        result |= PI_MEMORY_SCOPE_DEVICE | PI_MEMORY_SCOPE_SYSTEM;
+      }
+    }
+    if (paramValue) {
+      if (paramValueSize < sizeof(cl_device_atomic_capabilities))
+        return PI_ERROR_INVALID_VALUE;
+
+      std::memcpy(paramValue, &result, sizeof(result));
+      if (paramValueSizeRet)
+        *paramValueSizeRet = sizeof(result);
+    }
+    return PI_SUCCESS;
+  }
   case PI_DEVICE_INFO_ATOMIC_64: {
     cl_int ret_err = CL_SUCCESS;
     cl_bool result = CL_FALSE;

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -475,6 +475,10 @@ inline pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
        (ur_device_info_t)UR_DEVICE_INFO_BFLOAT16},
       {PI_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES,
        (ur_device_info_t)UR_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES},
+      {PI_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES,
+       (ur_device_info_t)UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES},
+      {PI_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES,
+       (ur_device_info_t)UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES},
   };
 
   auto InfoType = InfoMapping.find(ParamName);

sycl/plugins/unified_runtime/ur/adapters/level_zero/ur_level_zero.cpp

@@ -5,6 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===-----------------------------------------------------------------===//
+#include <sycl/detail/pi.h>
 
 #include <algorithm>
 #include <climits>
@@ -1164,6 +1165,26 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(
     // bfloat16 math functions are not yet supported on Intel GPUs.
     return ReturnValue(bool{false});
   }
+  case UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: {
+    // There are no explicit restrictions in L0 programming guide, so assume all
+    // are supported
+    pi_memory_order_capabilities result =
+        PI_MEMORY_ORDER_RELAXED | PI_MEMORY_ORDER_ACQUIRE |
+        PI_MEMORY_ORDER_RELEASE | PI_MEMORY_ORDER_ACQ_REL |
+        PI_MEMORY_ORDER_SEQ_CST;
+
+    return ReturnValue(result);
+  }
+  case UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES: {
+    // There are no explicit restrictions in L0 programming guide, so assume all
+    // are supported
+    pi_memory_scope_capabilities result =
+        PI_MEMORY_SCOPE_WORK_ITEM | PI_MEMORY_SCOPE_SUB_GROUP |
+        PI_MEMORY_SCOPE_WORK_GROUP | PI_MEMORY_SCOPE_DEVICE |
+        PI_MEMORY_SCOPE_SYSTEM;
+
+    return ReturnValue(result);
+  }
 
   // TODO: Implement.
   case UR_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES:

sycl/plugins/unified_runtime/ur/ur.hpp

@@ -43,6 +43,10 @@ const int UR_EXT_DEVICE_INFO_FREE_MEMORY = UR_EXT_DEVICE_INFO_END - 13;
 // const int ZER_EXT_DEVICE_INFO_DEVICE_ID = UR_EXT_DEVICE_INFO_END - 14;
 // const int ZER_EXT_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE =
 //     UR_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE;
+const int UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES =
+    UR_EXT_DEVICE_INFO_END - 16;
+const int UR_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES =
+    UR_EXT_DEVICE_INFO_END - 17;
 
 const ur_device_info_t UR_EXT_DEVICE_INFO_OPENCL_C_VERSION =
     (ur_device_info_t)0x103D;

sycl/source/detail/device_info.hpp

@@ -275,6 +275,19 @@ struct get_device_info_impl<std::vector<memory_order>,
   }
 };
 
+// Specialization for atomic_fence_order_capabilities, PI returns a bitfield
+template <>
+struct get_device_info_impl<std::vector<memory_order>,
+                            info::device::atomic_fence_order_capabilities> {
+  static std::vector<memory_order> get(RT::PiDevice dev, const plugin &Plugin) {
+    pi_memory_order_capabilities result;
+    Plugin.call_nocheck<PiApiKind::piDeviceGetInfo>(
+        dev, PiInfoCode<info::device::atomic_fence_order_capabilities>::value,
+        sizeof(pi_memory_order_capabilities), &result, nullptr);
+    return readMemoryOrderBitfield(result);
+  }
+};
+
 // Specialization for atomic_memory_scope_capabilities, PI returns a bitfield
 template <>
 struct get_device_info_impl<std::vector<memory_scope>,
@@ -288,6 +301,19 @@ struct get_device_info_impl<std::vector<memory_scope>,
   }
 };
 
+// Specialization for atomic_fence_scope_capabilities, PI returns a bitfield
+template <>
+struct get_device_info_impl<std::vector<memory_scope>,
+                            info::device::atomic_fence_scope_capabilities> {
+  static std::vector<memory_scope> get(RT::PiDevice dev, const plugin &Plugin) {
+    pi_memory_scope_capabilities result;
+    Plugin.call_nocheck<PiApiKind::piDeviceGetInfo>(
+        dev, PiInfoCode<info::device::atomic_fence_scope_capabilities>::value,
+        sizeof(pi_memory_scope_capabilities), &result, nullptr);
+    return readMemoryScopeBitfield(result);
+  }
+};
+
 // Specialization for bf16 math functions
 template <>
 struct get_device_info_impl<bool,
@@ -1005,13 +1031,27 @@ get_device_info_host<info::device::atomic_memory_order_capabilities>() {
           memory_order::acq_rel, memory_order::seq_cst};
 }
 
+template <>
+inline std::vector<memory_order>
+get_device_info_host<info::device::atomic_fence_order_capabilities>() {
+  return {memory_order::relaxed, memory_order::acquire, memory_order::release,
+          memory_order::acq_rel};
+}
+
 template <>
 inline std::vector<memory_scope>
 get_device_info_host<info::device::atomic_memory_scope_capabilities>() {
   return {memory_scope::work_item, memory_scope::sub_group,
           memory_scope::work_group, memory_scope::device, memory_scope::system};
 }
 
+template <>
+inline std::vector<memory_scope>
+get_device_info_host<info::device::atomic_fence_scope_capabilities>() {
+  return {memory_scope::work_item, memory_scope::sub_group,
+          memory_scope::work_group, memory_scope::device, memory_scope::system};
+}
+
 template <>
 inline bool
 get_device_info_host<info::device::ext_oneapi_bfloat16_math_functions>() {

sycl/test/abi/sycl_symbols_linux.dump

@@ -3,7 +3,7 @@
 # DO NOT EDIT IT MANUALLY. Refer to sycl/doc/developer/ABIPolicyGuide.md for more info.
 ################################################################################
 
-# RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so
+# RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir %python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so
 # REQUIRES: linux
 # UNSUPPORTED: libcxx
 
@@ -4332,6 +4332,8 @@ _ZNK4sycl3_V16device8get_infoINS0_4info6device29ext_oneapi_max_work_groups_2dEEE
 _ZNK4sycl3_V16device8get_infoINS0_4info6device29ext_oneapi_max_work_groups_3dEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device29preferred_vector_width_doubleEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device30partition_type_affinity_domainEEENS0_6detail19is_device_info_descIT_E11return_typeEv
+_ZNK4sycl3_V16device8get_infoINS0_4info6device31atomic_fence_order_capabilitiesEEENS0_6detail19is_device_info_descIT_E11return_typeEv
+_ZNK4sycl3_V16device8get_infoINS0_4info6device31atomic_fence_scope_capabilitiesEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device31ext_intel_gpu_hw_threads_per_euEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device32atomic_memory_order_capabilitiesEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device32atomic_memory_scope_capabilitiesEEENS0_6detail19is_device_info_descIT_E11return_typeEv