[SYCL] Initial support for sycl_ext_oneapi_atomic16 (#15158)

Spec:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_atomic16.asciidoc

This patch adds support for atomic_ref half specialization and
introduces a new atomic16 aspect.
There is no yet native support for 16bit float compare_exchange
function, extension is needed.

16bit int (incl. bfloat16) atomics will be supported later, when there
is a spec allowing these instructions.

Author: KornevNikita

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10614,7 +10614,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
               ",+SPV_INTEL_tensor_float32_conversion"
               ",+SPV_INTEL_optnone"
               ",+SPV_KHR_non_semantic_info"
-              ",+SPV_KHR_cooperative_matrix";
+              ",+SPV_KHR_cooperative_matrix"
+              ",+SPV_EXT_shader_atomic_float16_add";
   if (IsCPU)
     ExtArg += ",+SPV_INTEL_fp_max_error";
 

clang/test/Driver/sycl-spirv-ext-old-model.c

@@ -61,7 +61,8 @@
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_tensor_float32_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_optnone
 // CHECK-DEFAULT-SAME:,+SPV_KHR_non_semantic_info
-// CHECK-DEFAULT-SAME:,+SPV_KHR_cooperative_matrix"
+// CHECK-DEFAULT-SAME:,+SPV_KHR_cooperative_matrix
+// CHECK-DEFAULT-SAME:,+SPV_EXT_shader_atomic_float16_add"
 // CHECK-FPGA-HW: llvm-spirv{{.*}}"-spirv-ext=-all
 // CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_add
 // CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_min_max

clang/test/Driver/sycl-spirv-metadata-old-model.cpp

@@ -9,7 +9,7 @@
 // RUN:  FileCheck -check-prefix CHECK-WITHOUT %s
 
 // CHECK-WITH: llvm-spirv{{.*}} "--spirv-preserve-auxdata"
-// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_KHR_cooperative_matrix"
+// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_EXT_shader_atomic_float16_add"
 
 // CHECK-WITHOUT: "{{.*}}llvm-spirv"
 // CHECK-WITHOUT-NOT: --spirv-preserve-auxdata

clang/test/Driver/sycl-spirv-obj-old-model.cpp

@@ -11,7 +11,7 @@
 // SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]"
 // SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]"
 // SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata"
-// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_KHR_cooperative_matrix"
+// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_EXT_shader_atomic_float16_add"
 // SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]"
 // SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"
 // SPIRV_DEVICE_OBJ-SAME: "-fsycl-is-host"

clang/test/Driver/sycl-spirv-obj.cpp

@@ -11,7 +11,7 @@
 // SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]"
 // SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]"
 // SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata"
-// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_KHR_cooperative_matrix"
+// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_EXT_shader_atomic_float16_add"
 // SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]"
 // SPIRV_DEVICE_OBJ: clang-offload-packager{{.*}} "--image=file=[[DEVICE_SPV]]{{.*}}"
 // SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -84,6 +84,7 @@ def AspectExt_oneapi_queue_profiling_tag : Aspect<"ext_oneapi_queue_profiling_ta
 def AspectExt_oneapi_virtual_mem : Aspect<"ext_oneapi_virtual_mem">;
 def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group">;
 def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
+def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -146,7 +147,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
     AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_private_alloca, 
     AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group, 
-    AspectExt_intel_fpga_task_sequence],
+    AspectExt_intel_fpga_task_sequence,
+    AspectExt_oneapi_atomic16],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -505,6 +505,7 @@ extern __DPCPP_SYCL_EXTERNAL
   macro(__attribute__((opencl_global)), Arg)                                   \
       macro(__attribute__((opencl_local)), Arg) macro(, Arg)
 
+__SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, _Float16)
 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float)
 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double)
 __SPIRV_ATOMICS(__SPIRV_ATOMIC_SIGNED, int)

sycl/include/sycl/atomic_ref.hpp

@@ -25,6 +25,13 @@
 
 namespace sycl {
 inline namespace _V1 {
+
+// forward declarartion
+namespace detail::half_impl {
+class half;
+}
+using half = detail::half_impl::half;
+
 namespace detail {
 
 using memory_order = sycl::memory_order;
@@ -36,7 +43,7 @@ template <typename T> struct IsValidAtomicRefType {
        std::is_same_v<T, long> || std::is_same_v<T, unsigned long> ||
        std::is_same_v<T, long long> || std::is_same_v<T, unsigned long long> ||
        std::is_same_v<T, float> || std::is_same_v<T, double> ||
-       std::is_pointer_v<T>);
+       std::is_pointer_v<T> || std::is_same_v<T, sycl::half>);
 };
 
 template <sycl::access::address_space AS> struct IsValidAtomicRefAddressSpace {
@@ -116,7 +123,7 @@ class atomic_ref_base {
   static_assert(
       detail::IsValidAtomicRefType<T>::value,
       "Invalid atomic type.  Valid types are int, unsigned int, long, "
-      "unsigned long, long long, unsigned long long, float, double "
+      "unsigned long, long long, unsigned long long, sycl::half, float, double "
       "and pointer types");
   static_assert(detail::IsValidAtomicRefAddressSpace<AddressSpace>::value,
                 "Invalid atomic address_space.  Valid address spaces are: "
@@ -259,7 +266,7 @@ class atomic_ref_base {
 };
 
 // Hook allowing partial specializations to inherit atomic_ref_base
-template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder,
+template <typename T, size_t SizeOfT, memory_order DefaultOrder,
           memory_scope DefaultScope, access::address_space AddressSpace,
           typename = void>
 class atomic_ref_impl
@@ -270,10 +277,9 @@ class atomic_ref_impl
 };
 
 // Partial specialization for integral types
-template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder,
+template <typename T, size_t SizeOfT, memory_order DefaultOrder,
           memory_scope DefaultScope, access::address_space AddressSpace>
-class atomic_ref_impl<T, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope,
-                      AddressSpace,
+class atomic_ref_impl<T, SizeOfT, DefaultOrder, DefaultScope, AddressSpace,
                       typename std::enable_if_t<std::is_integral_v<T>>>
     : public atomic_ref_base<T, DefaultOrder, DefaultScope, AddressSpace> {
 
@@ -418,11 +424,11 @@ class atomic_ref_impl<T, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope,
 };
 
 // Partial specialization for floating-point types
-template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder,
+template <typename T, size_t SizeOfT, memory_order DefaultOrder,
           memory_scope DefaultScope, access::address_space AddressSpace>
-class atomic_ref_impl<T, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope,
-                      AddressSpace,
-                      typename std::enable_if_t<std::is_floating_point_v<T>>>
+class atomic_ref_impl<T, SizeOfT, DefaultOrder, DefaultScope, AddressSpace,
+                      typename std::enable_if_t<std::is_floating_point_v<T> ||
+                                                std::is_same_v<T, sycl::half>>>
     : public atomic_ref_base<T, DefaultOrder, DefaultScope, AddressSpace> {
 
 public:
@@ -534,15 +540,15 @@ class atomic_ref_impl<
 #else
 class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
 #endif
-    T, /*IsAspectAtomic64AttrUsed = */ true, DefaultOrder, DefaultScope,
-    AddressSpace, typename std::enable_if_t<std::is_integral_v<T>>>
-    : public atomic_ref_impl<T, /*IsAspectAtomic64AttrUsed = */ false,
-                             DefaultOrder, DefaultScope, AddressSpace> {
+    T, /*SizeOfT = */ 8, DefaultOrder, DefaultScope, AddressSpace,
+    typename std::enable_if_t<std::is_integral_v<T>>>
+    : public atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                             AddressSpace> {
 public:
-  using atomic_ref_impl<T, /*IsAspectAtomic64AttrUsed = */ false, DefaultOrder,
-                        DefaultScope, AddressSpace>::atomic_ref_impl;
-  using atomic_ref_impl<T, /*IsAspectAtomic64AttrUsed = */ false, DefaultOrder,
-                        DefaultScope, AddressSpace>::atomic_ref_impl::operator=;
+  using atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                        AddressSpace>::atomic_ref_impl;
+  using atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                        AddressSpace>::atomic_ref_impl::operator=;
 };
 
 // Partial specialization for 64-bit floating-point types needed for optional
@@ -554,28 +560,51 @@ class atomic_ref_impl<
 #else
 class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
 #endif
-    T, /*IsAspectAtomic64AttrUsed = */ true, DefaultOrder, DefaultScope,
-    AddressSpace, typename std::enable_if_t<std::is_floating_point_v<T>>>
-    : public atomic_ref_impl<T, /*IsAspectAtomic64AttrUsed = */ false,
-                             DefaultOrder, DefaultScope, AddressSpace> {
+    T, /*SizeOfT = */ 8, DefaultOrder, DefaultScope, AddressSpace,
+    typename std::enable_if_t<std::is_floating_point_v<T> ||
+                              std::is_same_v<T, sycl::half>>>
+    : public atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                             AddressSpace> {
 public:
-  using atomic_ref_impl<T, /*IsAspectAtomic64AttrUsed = */ false, DefaultOrder,
-                        DefaultScope, AddressSpace>::atomic_ref_impl;
-  using atomic_ref_impl<T, /*IsAspectAtomic64AttrUsed = */ false, DefaultOrder,
-                        DefaultScope, AddressSpace>::atomic_ref_impl::operator=;
+  using atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                        AddressSpace>::atomic_ref_impl;
+  using atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                        AddressSpace>::atomic_ref_impl::operator=;
+};
+
+// Partial specialization for 16-bit floating-point types needed for optional
+// kernel features
+template <typename T, memory_order DefaultOrder, memory_scope DefaultScope,
+          access::address_space AddressSpace>
+#ifndef __SYCL_DEVICE_ONLY__
+class atomic_ref_impl<
+#else
+class
+    [[__sycl_detail__::__uses_aspects__(aspect::ext_oneapi_atomic16)]] atomic_ref_impl<
+#endif
+    T, /*SizeOfT = */ 2, DefaultOrder, DefaultScope, AddressSpace,
+    typename std::enable_if_t<std::is_floating_point_v<T> ||
+                              std::is_same_v<T, sycl::half>>>
+    : public atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                             AddressSpace> {
+public:
+  using atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                        AddressSpace>::atomic_ref_impl;
+  using atomic_ref_impl<T, /*SizeOfT = */ 4, DefaultOrder, DefaultScope,
+                        AddressSpace>::atomic_ref_impl::operator=;
 };
 
 // Partial specialization for pointer types
 // Arithmetic is emulated because target's representation of T* is unknown
 // TODO: Find a way to use intptr_t or uintptr_t atomics instead
-template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder,
+template <typename T, size_t SizeOfT, memory_order DefaultOrder,
           memory_scope DefaultScope, access::address_space AddressSpace>
 #ifndef __SYCL_DEVICE_ONLY__
 class atomic_ref_impl<
 #else
 class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
 #endif
-    T *, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope, AddressSpace>
+    T *, SizeOfT, DefaultOrder, DefaultScope, AddressSpace>
     : public atomic_ref_base<uintptr_t, DefaultOrder, DefaultScope,
                              AddressSpace> {
 
@@ -713,15 +742,17 @@ template <typename T, memory_order DefaultOrder, memory_scope DefaultScope,
           access::address_space AddressSpace =
               access::address_space::generic_space>
 // if sizeof(T) == 8 bytes, then the type T is optional kernel feature, so it
-// was decorated with [[__sycl_detail__::__uses_aspects__(aspect::atomic64))]]
+// was decorated with [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]]
 // attribute in detail::atomic_ref_impl partial specializations above
-class atomic_ref
-    : public detail::atomic_ref_impl<T, sizeof(T) == 8, DefaultOrder,
-                                     DefaultScope, AddressSpace> {
+//
+// if sizeof(T) == 2 bytes, then decorated with
+// [[__sycl_detail__::__uses_aspects__(aspect::ext_oneapi_atomic16)]]
+class atomic_ref : public detail::atomic_ref_impl<T, sizeof(T), DefaultOrder,
+                                                  DefaultScope, AddressSpace> {
 public:
-  using detail::atomic_ref_impl<T, sizeof(T) == 8, DefaultOrder, DefaultScope,
+  using detail::atomic_ref_impl<T, sizeof(T), DefaultOrder, DefaultScope,
                                 AddressSpace>::atomic_ref_impl;
-  using detail::atomic_ref_impl<T, sizeof(T) == 8, DefaultOrder, DefaultScope,
+  using detail::atomic_ref_impl<T, sizeof(T), DefaultOrder, DefaultScope,
                                 AddressSpace>::operator=;
 };
 

sycl/include/sycl/detail/spirv.hpp

@@ -560,6 +560,17 @@ AtomicLoad(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
   return sycl::bit_cast<T>(ResultInt);
 }
 
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated>
+inline typename std::enable_if_t<std::is_same_v<T, half>, T>
+AtomicLoad(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
+           memory_order Order) {
+  auto *Ptr = GetMultiPtrDecoratedAs<_Float16>(MPtr);
+  auto SPIRVOrder = getMemorySemanticsMask(Order);
+  auto SPIRVScope = getScope(Scope);
+  return __spirv_AtomicLoad(Ptr, SPIRVScope, SPIRVOrder);
+}
+
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated>
 inline typename std::enable_if_t<std::is_integral<T>::value>
@@ -584,6 +595,17 @@ AtomicStore(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
   __spirv_AtomicStore(PtrInt, SPIRVScope, SPIRVOrder, ValueInt);
 }
 
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated>
+inline typename std::enable_if_t<std::is_same_v<T, half>>
+AtomicStore(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
+            memory_order Order, T Value) {
+  auto *Ptr = GetMultiPtrDecoratedAs<_Float16>(MPtr);
+  auto SPIRVOrder = getMemorySemanticsMask(Order);
+  auto SPIRVScope = getScope(Scope);
+  __spirv_AtomicStore(Ptr, SPIRVScope, SPIRVOrder, Value);
+}
+
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated>
 inline typename std::enable_if_t<std::is_integral<T>::value, T>
@@ -610,6 +632,17 @@ AtomicExchange(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
   return sycl::bit_cast<T>(ResultInt);
 }
 
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated>
+inline typename std::enable_if_t<std::is_same_v<half, T>, T>
+AtomicExchange(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
+               memory_order Order, T Value) {
+  auto *Ptr = GetMultiPtrDecoratedAs<_Float16>(MPtr);
+  auto SPIRVOrder = getMemorySemanticsMask(Order);
+  auto SPIRVScope = getScope(Scope);
+  return __spirv_AtomicExchange(Ptr, SPIRVScope, SPIRVOrder, Value);
+}
+
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated>
 inline typename std::enable_if_t<std::is_integral<T>::value, T>
@@ -643,6 +676,17 @@ AtomicFAdd(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
   return __spirv_AtomicFAddEXT(Ptr, SPIRVScope, SPIRVOrder, Value);
 }
 
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated>
+inline typename std::enable_if_t<std::is_same_v<half, T>, T>
+AtomicFAdd(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
+           memory_order Order, T Value) {
+  auto *Ptr = GetMultiPtrDecoratedAs<_Float16>(MPtr);
+  auto SPIRVOrder = getMemorySemanticsMask(Order);
+  auto SPIRVScope = getScope(Scope);
+  return __spirv_AtomicFAddEXT(Ptr, SPIRVScope, SPIRVOrder, Value);
+}
+
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated>
 inline typename std::enable_if_t<std::is_integral<T>::value, T>
@@ -698,6 +742,17 @@ AtomicMin(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
   return __spirv_AtomicMin(Ptr, SPIRVScope, SPIRVOrder, Value);
 }
 
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated>
+inline typename std::enable_if_t<std::is_same_v<half, T>, T>
+AtomicMin(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
+  auto *Ptr = GetMultiPtrDecoratedAs<_Float16>(MPtr);
+  auto SPIRVOrder = getMemorySemanticsMask(Order);
+  auto SPIRVScope = getScope(Scope);
+  return __spirv_AtomicFMinEXT(Ptr, SPIRVScope, SPIRVOrder, Value);
+}
+
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated>
 inline typename std::enable_if_t<std::is_integral<T>::value, T>
@@ -720,6 +775,17 @@ AtomicMax(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
   return __spirv_AtomicMax(Ptr, SPIRVScope, SPIRVOrder, Value);
 }
 
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated>
+inline typename std::enable_if_t<std::is_same_v<half, T>, T>
+AtomicMax(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
+          memory_order Order, T Value) {
+  auto *Ptr = GetMultiPtrDecoratedAs<_Float16>(MPtr);
+  auto SPIRVOrder = getMemorySemanticsMask(Order);
+  auto SPIRVScope = getScope(Scope);
+  return __spirv_AtomicFMaxEXT(Ptr, SPIRVScope, SPIRVOrder, Value);
+}
+
 // Native shuffles map directly to a shuffle intrinsic:
 // - The Intel SPIR-V extension natively supports all arithmetic types.
 //   However, OpenCL extension natively supports float vectors,

sycl/include/sycl/detail/type_traits.hpp

@@ -333,6 +333,9 @@ template <typename T>
 struct is_floating_point
     : is_floating_point_impl<std::remove_cv_t<get_elem_type_t<T>>> {};
 
+template <typename T>
+constexpr bool is_floating_point_v = is_floating_point<T>::value;
+
 // is_arithmetic
 template <typename T>
 struct is_arithmetic

sycl/include/sycl/device_aspect_macros.hpp

@@ -390,6 +390,11 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_sample_2d_usm__ 0
 #endif
 
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_atomic16__
+//__SYCL_ASPECT(ext_oneapi_oneapi_atomic16, 80)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_atomic16__ 0
+#endif
+
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
 // __SYCL_ASPECT(host, 0)
 #define __SYCL_ANY_DEVICE_HAS_host__ 0
@@ -769,3 +774,8 @@
 //__SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_sample_2d_usm__ 0
 #endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_atomic16__
+//__SYCL_ASPECT(ext_oneapi_oneapi_atomic16, 80)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_atomic16__ 0
+#endif

sycl/include/sycl/info/aspects.def

@@ -70,3 +70,4 @@ __SYCL_ASPECT(ext_oneapi_image_array, 76)
 __SYCL_ASPECT(ext_oneapi_unique_addressing_per_dim, 77)
 __SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 78)
 __SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79)
+__SYCL_ASPECT(ext_oneapi_atomic16, 80)