DPC++ Runtime
Runtime libraries for oneAPI DPC++
common.hpp
Go to the documentation of this file.
1 //==---------------- common.hpp - DPC++ Explicit SIMD API ----------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 // definitions used in Explicit SIMD APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
16 
17 #include <sycl/detail/defines.hpp>
18 
19 #include <cstdint> // for uint* types
20 #include <type_traits>
21 
23 
24 #ifdef __SYCL_DEVICE_ONLY__
25 #define __ESIMD_UNSUPPORTED_ON_HOST
26 #else // __SYCL_DEVICE_ONLY__
27 #define __ESIMD_UNSUPPORTED_ON_HOST \
28  throw sycl::exception(sycl::errc::feature_not_supported, \
29  "This ESIMD feature is not supported on HOST")
30 #endif // __SYCL_DEVICE_ONLY__
31 
33 
34 namespace sycl {
36 namespace ext::intel::esimd {
37 
40 
41 using uchar = unsigned char;
42 using ushort = unsigned short;
43 using uint = unsigned int;
44 
47 struct saturation_on_tag : std::true_type {};
48 
50 struct saturation_off_tag : std::false_type {};
51 
53 static inline constexpr saturation_off_tag saturation_off{};
54 
56 static inline constexpr saturation_on_tag saturation_on{};
57 
59 enum class rgba_channel : uint8_t { R, G, B, A };
60 
64 using SurfaceIndex = unsigned int;
65 
66 namespace detail {
67 
68 template <typename T>
70  static constexpr bool value =
71  std::is_same_v<T, __ESIMD_NS::saturation_on_tag> ||
72  std::is_same_v<T, __ESIMD_NS::saturation_off_tag>;
73 };
74 
75 template <class T>
77 
79 ESIMD_INLINE constexpr bool isPowerOf2(unsigned int n) {
80  return (n & (n - 1)) == 0;
81 }
82 
86 ESIMD_INLINE constexpr bool isPowerOf2(unsigned int n, unsigned int limit) {
87  return (n & (n - 1)) == 0 && n <= limit;
88 }
89 
90 template <rgba_channel Ch>
91 static inline constexpr uint8_t ch = 1 << static_cast<int>(Ch);
92 static inline constexpr uint8_t chR = ch<rgba_channel::R>;
93 static inline constexpr uint8_t chG = ch<rgba_channel::G>;
94 static inline constexpr uint8_t chB = ch<rgba_channel::B>;
95 static inline constexpr uint8_t chA = ch<rgba_channel::A>;
96 
97 // Shared Local Memory Binding Table Index (aka surface index).
98 static inline constexpr SurfaceIndex SLM_BTI = 254;
99 static inline constexpr SurfaceIndex INVALID_BTI =
100  static_cast<SurfaceIndex>(-1);
101 } // namespace detail
102 
105 enum class rgba_channel_mask : uint8_t {
106  R = detail::chR,
107  G = detail::chG,
109  B = detail::chB,
113  A = detail::chA,
121 };
122 
124  int Pos = static_cast<int>(Ch);
125  return (static_cast<int>(M) & (1 << Pos)) >> Pos;
126 }
127 
129  return is_channel_enabled(M, rgba_channel::R) +
130  is_channel_enabled(M, rgba_channel::G) +
131  is_channel_enabled(M, rgba_channel::B) +
132  is_channel_enabled(M, rgba_channel::A);
133 }
134 
135 #define __ESIMD_USM_DWORD_ATOMIC_TO_LSC \
136  " is supported only on ACM, PVC. USM-based atomic will be auto-converted " \
137  "to LSC version."
138 
145 enum class atomic_op : uint8_t {
147  add = 0x0,
149  sub = 0x1,
151  inc = 0x2,
153  dec = 0x3,
155  umin = 0x4,
156  min __SYCL_DEPRECATED("use umin") = umin,
158  umax = 0x5,
159  max __SYCL_DEPRECATED("use smax") = umax,
161  xchg = 0x6,
163  cmpxchg = 0x7,
165  bit_and = 0x8,
167  bit_or = 0x9,
169  bit_xor = 0xa,
171  smin = 0xb,
172  minsint __SYCL_DEPRECATED("use smin") = smin,
174  smax = 0xc,
175  maxsint __SYCL_DEPRECATED("use smax") = 0xc,
182  fcmpxchg = 0x12,
186  load = 0x15,
187  store = 0x16,
190  predec = 0xff,
191 };
192 
193 #undef __ESIMD_USM_DWORD_TO_LSC_MSG
194 
196 
197 namespace detail {
198 template <__ESIMD_NS::native::lsc::atomic_op Op> constexpr int get_num_args() {
199  if constexpr (Op == __ESIMD_NS::native::lsc::atomic_op::inc ||
201  Op == __ESIMD_NS::native::lsc::atomic_op::load) {
202  return 0;
203  } else if constexpr (Op == __ESIMD_NS::native::lsc::atomic_op::store ||
205  Op == __ESIMD_NS::native::lsc::atomic_op::sub ||
206  Op == __ESIMD_NS::native::lsc::atomic_op::smin ||
207  Op == __ESIMD_NS::native::lsc::atomic_op::smax ||
208  Op == __ESIMD_NS::native::lsc::atomic_op::umin ||
209  Op == __ESIMD_NS::native::lsc::atomic_op::umax ||
210  Op == __ESIMD_NS::native::lsc::atomic_op::fadd ||
211  Op == __ESIMD_NS::native::lsc::atomic_op::fsub ||
217  return 1;
218  } else if constexpr (Op == __ESIMD_NS::native::lsc::atomic_op::cmpxchg ||
219  Op == __ESIMD_NS::native::lsc::atomic_op::fcmpxchg) {
220  return 2;
221  } else {
222  return -1; // error
223  }
224 }
225 
226 template <__ESIMD_NS::atomic_op Op> constexpr bool has_lsc_equivalent() {
227  switch (Op) {
228  case __ESIMD_NS::atomic_op::xchg:
229  case __ESIMD_NS::atomic_op::predec:
230  return false;
231  default:
232  return true;
233  }
234 }
235 
236 template <__ESIMD_NS::atomic_op Op>
238  switch (Op) {
241  case __ESIMD_NS::atomic_op::sub:
242  return __ESIMD_NS::native::lsc::atomic_op::sub;
243  case __ESIMD_NS::atomic_op::inc:
244  return __ESIMD_NS::native::lsc::atomic_op::inc;
248  return __ESIMD_NS::native::lsc::atomic_op::umin;
250  return __ESIMD_NS::native::lsc::atomic_op::umax;
251  case __ESIMD_NS::atomic_op::cmpxchg:
252  return __ESIMD_NS::native::lsc::atomic_op::cmpxchg;
259  case __ESIMD_NS::atomic_op::minsint:
260  return __ESIMD_NS::native::lsc::atomic_op::smin;
261  case __ESIMD_NS::atomic_op::maxsint:
262  return __ESIMD_NS::native::lsc::atomic_op::smax;
267  case __ESIMD_NS::atomic_op::fcmpwr:
268  return __ESIMD_NS::native::lsc::atomic_op::fcmpxchg;
269  case __ESIMD_NS::atomic_op::fadd:
270  return __ESIMD_NS::native::lsc::atomic_op::fadd;
271  case __ESIMD_NS::atomic_op::fsub:
272  return __ESIMD_NS::native::lsc::atomic_op::fsub;
273  case __ESIMD_NS::atomic_op::load:
274  return __ESIMD_NS::native::lsc::atomic_op::load;
275  case __ESIMD_NS::atomic_op::store:
276  return __ESIMD_NS::native::lsc::atomic_op::store;
277  default:
278  static_assert(has_lsc_equivalent<Op>() && "Unsupported LSC atomic op");
279  }
280 }
281 
282 template <__ESIMD_NS::native::lsc::atomic_op Op>
284  switch (Op) {
287  case __ESIMD_NS::native::lsc::atomic_op::sub:
288  return __ESIMD_NS::atomic_op::sub;
289  case __ESIMD_NS::native::lsc::atomic_op::inc:
290  return __ESIMD_NS::atomic_op::inc;
293  case __ESIMD_NS::native::lsc::atomic_op::umin:
295  case __ESIMD_NS::native::lsc::atomic_op::umax:
297  case __ESIMD_NS::native::lsc::atomic_op::cmpxchg:
298  return __ESIMD_NS::atomic_op::cmpxchg;
305  case __ESIMD_NS::native::lsc::atomic_op::smin:
306  return __ESIMD_NS::atomic_op::minsint;
307  case __ESIMD_NS::native::lsc::atomic_op::smax:
308  return __ESIMD_NS::atomic_op::maxsint;
313  case __ESIMD_NS::native::lsc::atomic_op::fcmpxchg:
314  return __ESIMD_NS::atomic_op::fcmpwr;
315  case __ESIMD_NS::native::lsc::atomic_op::fadd:
316  return __ESIMD_NS::atomic_op::fadd;
317  case __ESIMD_NS::native::lsc::atomic_op::fsub:
318  return __ESIMD_NS::atomic_op::fsub;
319  case __ESIMD_NS::native::lsc::atomic_op::load:
320  return __ESIMD_NS::atomic_op::load;
321  case __ESIMD_NS::native::lsc::atomic_op::store:
322  return __ESIMD_NS::atomic_op::store;
323  }
324 }
325 
326 template <__ESIMD_NS::atomic_op Op> constexpr int get_num_args() {
327  if constexpr (has_lsc_equivalent<Op>()) {
328  return get_num_args<to_lsc_atomic_op<Op>()>();
329  } else {
330  switch (Op) {
331  case __ESIMD_NS::atomic_op::xchg:
332  case __ESIMD_NS::atomic_op::predec:
333  return 1;
334  default:
335  return -1; // error
336  }
337  }
338 }
339 
340 } // namespace detail
341 
342 } // namespace ext::intel::esimd
343 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
344 } // namespace sycl
sycl::_V1::ext::intel::esimd::rgba_channel_mask
rgba_channel_mask
Represents a pixel's channel mask - all possible combinations of enabled channels.
Definition: common.hpp:105
sycl::_V1::ext::intel::esimd::atomic_op::sub
@ sub
Subtraction: *addr = *addr - src0.
common.hpp
__ESIMD_USM_DWORD_ATOMIC_TO_LSC
#define __ESIMD_USM_DWORD_ATOMIC_TO_LSC
Definition: common.hpp:135
sycl::_V1::ext::oneapi::bit_and
std::bit_and< T > bit_and
Definition: functional.hpp:24
sycl::_V1::ext::intel::esimd::rgba_channel_mask::AGR
@ AGR
sycl::_V1::ext::intel::esimd::atomic_op::inc
@ inc
Increment: *addr = *addr + 1.
sycl::_V1::ext::intel::esimd::rgba_channel_mask::GR
@ GR
sycl::_V1::ext::intel::esimd::detail::chB
static constexpr uint8_t chB
Definition: common.hpp:94
sycl::_V1::ext::intel::esimd::atomic_op::xchg
@ xchg
Exchange. *addr == src0;
sycl::_V1::ext::intel::esimd::detail::get_num_args
constexpr int get_num_args()
Definition: common.hpp:198
sycl::_V1::ext::intel::esimd::detail::is_saturation_tag_v
constexpr bool is_saturation_tag_v
Definition: common.hpp:76
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::ext::intel::esimd::atomic_op::umin
@ umin
Minimum: *addr = min(*addr, src0).
sycl::_V1::ext::intel::esimd::rgba_channel_mask::AB
@ AB
sycl::_V1::ext::oneapi::bit_xor
std::bit_xor< T > bit_xor
Definition: functional.hpp:23
defines_elementary.hpp
sycl::_V1::ext::intel::esimd::atomic_op::cmpxchg
@ cmpxchg
Compare and exchange. if (*addr == src0) *sddr = src1;
sycl::_V1::bit_and
std::bit_and< T > bit_and
Definition: functional.hpp:20
sycl::_V1::ext::intel::esimd::atomic_op::load
@ load
sycl::_V1::bit_or
std::bit_or< T > bit_or
Definition: functional.hpp:21
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::ext::intel::esimd::rgba_channel::G
@ G
sycl::_V1::ext::intel::esimd::detail::to_lsc_atomic_op
constexpr sycl::ext::intel::esimd::native::lsc::atomic_op to_lsc_atomic_op()
Definition: common.hpp:237
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::_V1::ext::intel::esimd::uchar
unsigned char uchar
Definition: common.hpp:41
sycl::_V1::ext::intel::esimd::rgba_channel_mask::BR
@ BR
sycl::_V1::ext::intel::esimd::rgba_channel
rgba_channel
Represents a pixel's channel.
Definition: common.hpp:59
sycl::_V1::ext::intel::esimd::detail::isPowerOf2
constexpr ESIMD_INLINE bool isPowerOf2(unsigned int n, unsigned int limit)
Check at compile time if given 32 bit positive integer is both:
Definition: common.hpp:86
sycl::_V1::ext::intel::esimd::detail::has_lsc_equivalent
constexpr bool has_lsc_equivalent()
Definition: common.hpp:226
sycl::_V1::ext::oneapi::fmax
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmax(T x, T y)
Definition: bf16_storage_builtins.hpp:60
sycl::_V1::ext::intel::esimd::saturation_on_tag
Gen hardware supports applying saturation to results of certain operations.
Definition: common.hpp:47
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:46
sycl::_V1::ext::intel::esimd::rgba_channel::R
@ R
sycl::_V1::ext::intel::esimd::rgba_channel_mask::AR
@ AR
sycl::_V1::ext::intel::esimd::atomic_op::predec
@ predec
Decrement: *addr = *addr - 1.
sycl::_V1::ext::intel::esimd::SurfaceIndex
unsigned int SurfaceIndex
Surface index type.
Definition: common.hpp:64
sycl::_V1::bit_xor
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
sycl::_V1::ext::intel::esimd::saturation_off_tag
This type tag represents "saturation off" behavior.
Definition: common.hpp:50
sycl::_V1::ext::oneapi::fmin
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmin(T x, T y)
Definition: bf16_storage_builtins.hpp:49
sycl::_V1::ext::intel::esimd::saturation_off
static constexpr saturation_off_tag saturation_off
Type tag object representing "saturation off" behavior.
Definition: common.hpp:53
sycl::_V1::ext::intel::esimd::detail::chA
static constexpr uint8_t chA
Definition: common.hpp:95
defines.hpp
sycl::_V1::ext::intel::esimd::rgba_channel::A
@ A
sycl::_V1::ext::intel::esimd::atomic_op::smin
@ smin
Minimum (signed integer): *addr = min(*addr, src0).
sycl::_V1::ext::intel::esimd::detail::SLM_BTI
static constexpr SurfaceIndex SLM_BTI
Definition: common.hpp:98
common.hpp
sycl::_V1::ext::intel::esimd::rgba_channel_mask::ABG
@ ABG
sycl::_V1::ext::intel::esimd::rgba_channel::B
@ B
sycl::_V1::ext::intel::esimd::detail::INVALID_BTI
static constexpr SurfaceIndex INVALID_BTI
Definition: common.hpp:99
sycl::_V1::dec
constexpr stream_manipulator dec
Definition: stream.hpp:744
sycl::_V1::ext::intel::esimd::atomic_op::umax
@ umax
Maximum: *addr = max(*addr, src0).
sycl::_V1::ext::intel::esimd::atomic_op::fcmpxchg
@ fcmpxchg
Compare and exchange (floating point).
sycl::_V1::ext::intel::esimd::saturation_on
static constexpr saturation_on_tag saturation_on
Type tag object representing "saturation on" behavior.
Definition: common.hpp:56
sycl::_V1::ext::intel::esimd::ushort
unsigned short ushort
Definition: common.hpp:42
sycl::_V1::ext::intel::esimd::detail::chR
static constexpr uint8_t chR
Definition: common.hpp:92
sycl::_V1::ext::intel::esimd::uint
unsigned int uint
Definition: common.hpp:43
sycl::_V1::ext::intel::esimd::rgba_channel_mask::ABR
@ ABR
sycl::_V1::ext::intel::esimd::atomic_op::store
@ store
sycl::_V1::ext::intel::esimd::detail::chG
static constexpr uint8_t chG
Definition: common.hpp:93
sycl::_V1::ext::intel::esimd::native::lsc::atomic_op
atomic_op
LSC atomic operation codes.
Definition: common.hpp:39
sycl::_V1::ext::intel::esimd::detail::is_saturation_tag
Definition: common.hpp:69
sycl::_V1::ext::intel::esimd::is_channel_enabled
constexpr int is_channel_enabled(rgba_channel_mask M, rgba_channel Ch)
Definition: common.hpp:123
sycl::_V1::ext::intel::esimd::rgba_channel_mask::BG
@ BG
sycl::_V1::ext::intel::esimd::rgba_channel_mask::BGR
@ BGR
sycl::_V1::ext::intel::esimd::detail::to_atomic_op
constexpr sycl::ext::intel::esimd::atomic_op to_atomic_op()
Definition: common.hpp:283
sycl::_V1::ext::intel::esimd::rgba_channel_mask::ABGR
@ ABGR
sycl::_V1::detail::device_global_map::add
void add(const void *DeviceGlobalPtr, const char *UniqueId)
Definition: device_global_map.cpp:15
sycl::_V1::ext::intel::esimd::detail::ch
static constexpr uint8_t ch
Definition: common.hpp:91
sycl::_V1::ext::intel::esimd::get_num_channels_enabled
constexpr int get_num_channels_enabled(rgba_channel_mask M)
Definition: common.hpp:128
sycl::_V1::ext::intel::esimd::atomic_op::smax
@ smax
Maximum (signed integer): *addr = max(*addr, src0).
sycl::_V1::ext::intel::esimd::atomic_op
atomic_op
Represents an atomic operation.
Definition: common.hpp:145
sycl::_V1::ext::intel::esimd::rgba_channel_mask::AG
@ AG
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
sycl::_V1::ext::oneapi::bit_or
std::bit_or< T > bit_or
Definition: functional.hpp:22