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 experimental Explicit SIMD APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
14 
16 
17 // Macros for internal use
18 #define __ESIMD_ENS sycl::ext::intel::experimental::esimd
19 #define __ESIMD_EDNS sycl::ext::intel::experimental::esimd::detail
20 
22 
24 namespace __ESIMD_ENS {
25 
28 
29 enum class argument_type {
30  U1 = 0, // unsigned 1 bit
31  S1 = 1, // signed 1 bit
32  U2 = 2, // unsigned 2 bits
33  S2 = 3, // signed 2 bits
34  U4 = 4, // unsigned 4 bits
35  S4 = 5, // signed 4 bits
36  U8 = 6, // unsigned 8 bits
37  S8 = 7, // signed 8 bits
38  BF16 = 8, // bfloat 16
39  FP16 = 9, // half float
40  TF32 = 11 // tensorfloat 32
41 };
42 
45 enum class lsc_scope : uint8_t {
46  group = 0,
47  local = 1,
48  tile = 2,
49  gpu = 3,
50  gpus = 4,
51  system = 5,
52  sysacq = 6,
53 };
54 
57 enum class lsc_fence_op : uint8_t {
58  none = 0,
59  evict = 1,
60  invalidate = 2,
61  discard = 3,
62  clean = 4,
63  flushl3 = 5,
65 };
66 
69 enum class lsc_memory_kind : uint8_t {
70  untyped_global = 0,
72  typed_global = 2,
73  shared_local = 3,
74 };
75 
77 enum class lsc_data_size : uint8_t {
78  default_size = 0,
79  u8 = 1,
80  u16 = 2,
81  u32 = 3,
82  u64 = 4,
83  u8u32 = 5,
84  u16u32 = 6,
85  u16u32h = 7,
86 };
87 
88 namespace detail {
90 enum class lsc_atomic_op : uint8_t {
91  iinc = 0x08, // atomic integer increment
92  idec = 0x09, // atomic integer decrement
93  load = 0x0a, // atomic load
94  store = 0x0b, // atomic store
95  iadd = 0x0c, // atomic integer add
96  isub = 0x0d, // atomic integer subtract
97  smin = 0x0e, // atomic signed int min
98  smax = 0x0f, // atomic signed int max
99  umin = 0x10, // atomic unsigned int min
100  umax = 0x11, // atomic unsigned int max
101  icas = 0x12, // atomic int compare and swap
102  fadd = 0x13, // floating-point add
103  fsub = 0x14, // floating-point subtract
104  fmin = 0x15, // floating-point min
105  fmax = 0x16, // floating-point max
106  fcas = 0x17, // floating-point CAS
107  bit_and = 0x18, // logical (bitwise) AND
108  bit_or = 0x19, // logical (bitwise) OR
109  bit_xor = 0x1a, // logical (bitwise) XOR
110 };
111 
112 enum class lsc_vector_size : uint8_t {
113  n1 = 1,
114  n2 = 2,
115  n3 = 3,
116  n4 = 4,
117  n8 = 5,
118  n16 = 6,
119  n32 = 7,
120  n64 = 8,
121 };
122 
123 enum class lsc_data_order : uint8_t {
124  nontranspose = 1,
125  transpose = 2,
126 };
127 
128 template <lsc_vector_size VS> constexpr void check_lsc_vector_size() {
129  static_assert(VS == lsc_vector_size::n1 || VS == lsc_vector_size::n2 ||
130  VS == lsc_vector_size::n3 || VS == lsc_vector_size::n4 ||
131  VS == lsc_vector_size::n8 || VS == lsc_vector_size::n16 ||
132  VS == lsc_vector_size::n64 || VS == lsc_vector_size::n32,
133  "Unsupported vector size");
134 }
135 
136 template <uint8_t VS> constexpr void check_lsc_vector_size() {
137  static_assert(VS == 1 || VS == 2 || VS == 3 || VS == 4 || VS == 8 ||
138  VS == 16 || VS == 32 || VS == 64,
139  "Unsupported vector size");
140 }
141 
142 template <typename T, lsc_data_size DS> constexpr void check_lsc_data_size() {
143  static_assert(DS != lsc_data_size::default_size || sizeof(T) == 1 ||
144  sizeof(T) == 2 || sizeof(T) == 4 || sizeof(T) == 8,
145  "Unsupported data type");
146 }
147 
148 template <__ESIMD_NS::atomic_op Op> constexpr void check_lsc_atomic_op() {
149  static_assert(Op == __ESIMD_NS::atomic_op::add ||
150  Op == __ESIMD_NS::atomic_op::sub ||
151  Op == __ESIMD_NS::atomic_op::inc ||
155  Op == __ESIMD_NS::atomic_op::cmpxchg ||
159  Op == __ESIMD_NS::atomic_op::minsint ||
160  Op == __ESIMD_NS::atomic_op::maxsint ||
163  Op == __ESIMD_NS::atomic_op::fcmpwr ||
164  Op == __ESIMD_NS::atomic_op::fadd ||
165  Op == __ESIMD_NS::atomic_op::fsub ||
166  Op == __ESIMD_NS::atomic_op::load ||
167  Op == __ESIMD_NS::atomic_op::store,
168  "Unsupported operation for LSC atomics");
169 }
170 
172 template <__ESIMD_NS::atomic_op Op, unsigned NumSrc>
173 constexpr void check_lsc_atomic() {
174  check_lsc_atomic_op<Op>();
175  if constexpr (Op == __ESIMD_NS::atomic_op::inc ||
177  Op == __ESIMD_NS::atomic_op::load) {
178  static_assert(NumSrc == 0, "No source operands are expected");
179  }
180  if constexpr (Op == __ESIMD_NS::atomic_op::store ||
182  Op == __ESIMD_NS::atomic_op::sub ||
183  Op == __ESIMD_NS::atomic_op::minsint ||
184  Op == __ESIMD_NS::atomic_op::maxsint ||
187  Op == __ESIMD_NS::atomic_op::fadd ||
188  Op == __ESIMD_NS::atomic_op::fsub ||
194  static_assert(NumSrc == 1, "One source operand is expected");
195  }
196  if constexpr (Op == __ESIMD_NS::atomic_op::cmpxchg ||
197  Op == __ESIMD_NS::atomic_op::fcmpwr) {
198  static_assert(NumSrc == 2, "Two source operands are expected");
199  }
200 }
201 
202 template <__ESIMD_NS::atomic_op Op> constexpr lsc_atomic_op to_lsc_atomic_op() {
203  check_lsc_atomic_op<Op>();
204  switch (Op) {
206  return lsc_atomic_op::iadd;
207  case __ESIMD_NS::atomic_op::sub:
208  return lsc_atomic_op::isub;
209  case __ESIMD_NS::atomic_op::inc:
210  return lsc_atomic_op::iinc;
212  return lsc_atomic_op::idec;
214  return lsc_atomic_op::umin;
216  return lsc_atomic_op::umax;
217  case __ESIMD_NS::atomic_op::cmpxchg:
218  return lsc_atomic_op::icas;
220  return lsc_atomic_op::bit_and;
222  return lsc_atomic_op::bit_or;
224  return lsc_atomic_op::bit_xor;
225  case __ESIMD_NS::atomic_op::minsint:
226  return lsc_atomic_op::smin;
227  case __ESIMD_NS::atomic_op::maxsint:
228  return lsc_atomic_op::smax;
230  return lsc_atomic_op::fmax;
232  return lsc_atomic_op::fmin;
233  case __ESIMD_NS::atomic_op::fcmpwr:
234  return lsc_atomic_op::fcas;
235  case __ESIMD_NS::atomic_op::fadd:
236  return lsc_atomic_op::fadd;
237  case __ESIMD_NS::atomic_op::fsub:
238  return lsc_atomic_op::fsub;
239  case __ESIMD_NS::atomic_op::load:
240  return lsc_atomic_op::load;
241  case __ESIMD_NS::atomic_op::store:
242  return lsc_atomic_op::store;
243  default:
244  return lsc_atomic_op::iinc;
245  }
246 }
247 
248 template <lsc_vector_size VS> constexpr uint8_t to_int() {
249  check_lsc_vector_size<VS>();
250  switch (VS) {
251  case lsc_vector_size::n1:
252  return 1;
253  case lsc_vector_size::n2:
254  return 2;
255  case lsc_vector_size::n3:
256  return 3;
257  case lsc_vector_size::n4:
258  return 4;
259  case lsc_vector_size::n8:
260  return 8;
261  case lsc_vector_size::n16:
262  return 16;
263  case lsc_vector_size::n32:
264  return 32;
265  case lsc_vector_size::n64:
266  return 64;
267  default:
268  return 1;
269  }
270 }
271 
272 template <uint8_t VS> constexpr lsc_vector_size to_lsc_vector_size() {
273  check_lsc_vector_size<VS>();
274  switch (VS) {
275  case 1:
276  return lsc_vector_size::n1;
277  case 2:
278  return lsc_vector_size::n2;
279  case 3:
280  return lsc_vector_size::n3;
281  case 4:
282  return lsc_vector_size::n4;
283  case 8:
284  return lsc_vector_size::n8;
285  case 16:
286  return lsc_vector_size::n16;
287  case 32:
288  return lsc_vector_size::n32;
289  case 64:
290  return lsc_vector_size::n64;
291  default:
292  return lsc_vector_size::n1;
293  }
294 }
295 
296 template <typename T, lsc_data_size DS>
298  check_lsc_data_size<T, DS>();
299  if (DS != lsc_data_size::default_size)
300  return DS;
301  else if (sizeof(T) == 1)
302  return lsc_data_size::u8;
303  else if (sizeof(T) == 2)
304  return lsc_data_size::u16;
305  else if (sizeof(T) == 4)
306  return lsc_data_size::u32;
307  else if (sizeof(T) == 8)
308  return lsc_data_size::u64;
309  else
310  return DS;
311 }
312 
314  if (DS == lsc_data_size::u8)
315  return lsc_data_size::u8u32;
316  if (DS == lsc_data_size::u16)
317  return lsc_data_size::u16u32;
318  return DS;
319 }
320 
321 template <typename T> struct lsc_expand_type {
322  using type = typename std::conditional<sizeof(T) < 4, uint32_t, T>::type;
323 };
324 
325 template <typename T> struct lsc_bitcast_type {
326 private:
327  using _type1 = typename std::conditional<sizeof(T) == 2, uint16_t, T>::type;
328  using _type2 = typename std::conditional<sizeof(T) == 1, uint8_t, T>::type;
329 
330 public:
331  using type =
332  typename std::conditional<sizeof(_type2) == 1, _type2, _type1>::type;
333 };
334 
335 } // namespace detail
336 
338 enum class cache_hint : uint8_t {
339  none = 0,
340  uncached = 1,
341  cached = 2,
342  write_back = 3,
343  write_through = 4,
344  streaming = 5,
345  read_invalidate = 6
346 };
347 
348 namespace detail {
349 
350 template <cache_hint Hint> class cache_hint_wrap {
351  template <cache_hint...> class is_one_of_t;
352  template <cache_hint Last>
353  struct is_one_of_t<Last>
354  : std::conditional<Last == Hint, std::true_type, std::false_type>::type {
355  };
356  template <cache_hint Head, cache_hint... Tail>
357  struct is_one_of_t<Head, Tail...>
358  : std::conditional<Head == Hint, std::true_type,
359  is_one_of_t<Tail...>>::type {};
360 
361 public:
362  constexpr operator cache_hint() const { return Hint; }
363  template <cache_hint... Hints> constexpr bool is_one_of() const {
364  return is_one_of_t<Hints...>::value;
365  }
366 };
367 
368 constexpr bool are_both(cache_hint First, cache_hint Second, cache_hint Val) {
369  return First == Val && Second == Val;
370 }
371 
372 enum class lsc_action { prefetch, load, store, atomic };
373 
374 template <lsc_action Action, cache_hint L1, cache_hint L3>
375 constexpr void check_lsc_cache_hint() {
376  constexpr auto L1H = cache_hint_wrap<L1>{};
377  constexpr auto L3H = cache_hint_wrap<L3>{};
378  if constexpr (Action == lsc_action::prefetch) {
379  static_assert(
380  L1H.template is_one_of<cache_hint::cached, cache_hint::uncached,
381  cache_hint::streaming>() &&
382  L3H.template is_one_of<cache_hint::cached,
383  cache_hint::uncached>() &&
384  !are_both(L1H, L3H, cache_hint::uncached),
385  "unsupported cache hint");
386  } else if constexpr (Action == lsc_action::load) {
387  static_assert(
388  are_both(L1H, L3H, cache_hint::none) ||
389  (L1H.template is_one_of<cache_hint::uncached, cache_hint::cached,
390  cache_hint::streaming>() &&
391  L3H.template is_one_of<cache_hint::uncached,
392  cache_hint::cached>()) ||
393  (L1H == cache_hint::read_invalidate && L3H == cache_hint::cached),
394  "unsupported cache hint");
395  } else if constexpr (Action == lsc_action::store) {
396  static_assert(are_both(L1H, L3H, cache_hint::none) ||
397  are_both(L1H, L3H, cache_hint::write_back) ||
398  (L1H.template is_one_of<cache_hint::uncached,
399  cache_hint::write_through,
400  cache_hint::streaming>() &&
401  L3H.template is_one_of<cache_hint::uncached,
402  cache_hint::write_back>()),
403  "unsupported cache hint");
404  } else if constexpr (Action == lsc_action::atomic) {
405  static_assert(are_both(L1H, L3H, cache_hint::none) ||
406  (L1H == cache_hint::uncached &&
407  L3H.template is_one_of<cache_hint::uncached,
408  cache_hint::write_back>()),
409  "unsupported cache hint");
410  }
411 }
412 
413 } // namespace detail
414 
416 enum class split_barrier_action : uint8_t {
417  wait = 0, // split barrier wait
418  signal = 1, // split barrier signal
419 };
420 
422 
423 } // namespace __ESIMD_ENS
424 } // __SYCL_INLINE_NAMESPACE(cl)
cl::__ESIMD_ENS::detail::lsc_vector_size::n64
@ n64
cl::__ESIMD_ENS::detail::lsc_atomic_op::fcas
@ fcas
cl::__ESIMD_ENS::detail::lsc_vector_size::n4
@ n4
cl::__ESIMD_ENS::detail::lsc_bitcast_type
Definition: common.hpp:325
cl::__ESIMD_ENS::lsc_fence_op::discard
@ discard
invalidate all clean lines
cl::__ESIMD_ENS::detail::lsc_atomic_op::umax
@ umax
cl::__ESIMD_ENS::detail::lsc_atomic_op::umin
@ umin
cl::sycl::bit_xor
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
cl::__ESIMD_ENS::lsc_scope
lsc_scope
The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC.
Definition: common.hpp:45
cl::__ESIMD_ENS::lsc_data_size::u16u32
@ u16u32
load 8b, zero extend to 32b; store the opposite
cl::sycl::detail::device_global_map::add
void add(const void *DeviceGlobalPtr, const char *UniqueId)
Definition: device_global_map.cpp:16
cl::__ESIMD_ENS::detail::check_lsc_cache_hint
constexpr void check_lsc_cache_hint()
Definition: common.hpp:375
cl::__ESIMD_ENS::detail::lsc_vector_size::n1
@ n1
cl::__ESIMD_ENS::detail::lsc_atomic_op
lsc_atomic_op
LSC atomic operations op codes.
Definition: common.hpp:90
cl::__ESIMD_ENS::cache_hint::cached
@ cached
cl::__ESIMD_ENS::detail::lsc_atomic_op::icas
@ icas
cl::__ESIMD_ENS::detail::lsc_atomic_op::iadd
@ iadd
cl::__ESIMD_ENS::detail::lsc_atomic_op::smin
@ smin
cl::__ESIMD_ENS::detail::to_lsc_vector_size
constexpr lsc_vector_size to_lsc_vector_size()
Definition: common.hpp:272
cl::__ESIMD_ENS::argument_type::S2
@ S2
cl::__ESIMD_ENS::lsc_scope::sysacq
@ sysacq
the entire system memory space
cl::__ESIMD_ENS::lsc_data_size::default_size
@ default_size
cl::__ESIMD_ENS::lsc_data_size
lsc_data_size
Data size or format to read or store.
Definition: common.hpp:77
cl::__ESIMD_ENS::lsc_memory_kind::shared_local
@ shared_local
typed global memory
cl::__ESIMD_ENS::argument_type::S1
@ S1
cl::__ESIMD_ENS::detail::lsc_action
lsc_action
Definition: common.hpp:372
cl::__ESIMD_ENS::detail::are_both
constexpr bool are_both(cache_hint First, cache_hint Second, cache_hint Val)
Definition: common.hpp:368
cl::__ESIMD_ENS::argument_type::S4
@ S4
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
cl::__ESIMD_ENS::detail::to_int
constexpr uint8_t to_int()
Definition: common.hpp:248
cl::__ESIMD_ENS::argument_type::TF32
@ TF32
cl::__ESIMD_ENS::lsc_fence_op
lsc_fence_op
The lsc_fence operation to apply to caches Supported platforms: DG2, PVC.
Definition: common.hpp:57
cl::__ESIMD_ENS::argument_type::FP16
@ FP16
cl::__ESIMD_ENS::lsc_data_size::u16
@ u16
cl::__ESIMD_ENS::lsc_fence_op::clean
@ clean
direct and clean lines are discarded w/o eviction
cl::__ESIMD_ENS::argument_type::S8
@ S8
cl::__ESIMD_ENS::detail::check_lsc_vector_size
constexpr void check_lsc_vector_size()
Definition: common.hpp:128
cl::__ESIMD_ENS::detail::check_lsc_atomic
constexpr void check_lsc_atomic()
Check the legality of lsc xatomic call in terms of size and type.
Definition: common.hpp:173
cl::__ESIMD_ENS::detail::lsc_data_order
lsc_data_order
Definition: common.hpp:123
cl::__ESIMD_ENS::lsc_data_size::u16u32h
@ u16u32h
load 16b, zero extend to 32b; store the opposite
cl::__ESIMD_ENS::cache_hint::write_back
@ write_back
cl::sycl::bit_or
std::bit_or< T > bit_or
Definition: functional.hpp:21
cl::__ESIMD_ENS::lsc_data_size::u32
@ u32
cl::__ESIMD_ENS::detail::expand_data_size
constexpr lsc_data_size expand_data_size(lsc_data_size DS)
Definition: common.hpp:313
cl::__ESIMD_ENS::lsc_memory_kind::untyped_global_low_pri
@ untyped_global_low_pri
untyped global memory
cl::__ESIMD_ENS::lsc_memory_kind
lsc_memory_kind
The specific LSC shared function to fence with lsc_fence Supported platforms: DG2,...
Definition: common.hpp:69
cl::__ESIMD_ENS::cache_hint
cache_hint
L1 or L3 cache hint kinds.
Definition: common.hpp:338
cl::__ESIMD_ENS::lsc_memory_kind::untyped_global
@ untyped_global
cl::__ESIMD_ENS::detail::lsc_vector_size::n3
@ n3
cl::__ESIMD_ENS::lsc_data_size::u8
@ u8
cl::__ESIMD_ENS::lsc_scope::group
@ group
cl::__ESIMD_ENS::split_barrier_action
split_barrier_action
Represents a split barrier action.
Definition: common.hpp:416
cl::sycl::fmax
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmax(T x, T y) __NOEXC
Definition: builtins.hpp:203
cl::sycl::ext::intel::experimental::prefetch
prefetch_impl< _B > prefetch
Definition: fpga_lsu.hpp:47
cl::sycl::fmin
detail::enable_if_t< detail::is_genfloat< T >::value, T > fmin(T x, T y) __NOEXC
Definition: builtins.hpp:216
cl::__ESIMD_ENS::argument_type
argument_type
Definition: common.hpp:29
cl::__ESIMD_ENS::detail::cache_hint_wrap::is_one_of
constexpr bool is_one_of() const
Definition: common.hpp:363
cl::__ESIMD_ENS::lsc_data_size::u64
@ u64
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::bit_and
std::bit_and< T > bit_and
Definition: functional.hpp:20
cl::__ESIMD_ENS::detail::check_lsc_atomic_op
constexpr void check_lsc_atomic_op()
Definition: common.hpp:148
cl::__ESIMD_ENS::lsc_memory_kind::typed_global
@ typed_global
low-priority untyped global memory
cl::__ESIMD_ENS::detail::lsc_atomic_op::idec
@ idec
cl::__ESIMD_ENS::split_barrier_action::signal
@ signal
cl::__ESIMD_ENS::detail::lsc_bitcast_type::type
typename std::conditional< sizeof(_type2)==1, _type2, _type1 >::type type
Definition: common.hpp:332
cl::__ESIMD_ENS::detail::lsc_vector_size::n2
@ n2
cl::__ESIMD_ENS::lsc_fence_op::invalidate
@ invalidate
dirty lines evicted and invalidated from L1
cl::__ESIMD_ENS::detail::cache_hint_wrap
Definition: common.hpp:350
cl::__ESIMD_ENS::detail::lsc_atomic_op::smax
@ smax
cl::__ESIMD_ENS::cache_hint::write_through
@ write_through
cl::__ESIMD_ENS::lsc_scope::gpus
@ gpus
entire GPU, flush out to the GPUs LLC
cl::__ESIMD_ENS::argument_type::U2
@ U2
cl::__ESIMD_ENS::argument_type::U4
@ U4
cl::__ESIMD_ENS::cache_hint::read_invalidate
@ read_invalidate
cl::__ESIMD_ENS::detail::lsc_expand_type
Definition: common.hpp:321
cl::__ESIMD_ENS::detail::lsc_data_order::nontranspose
@ nontranspose
common.hpp
cl::sycl::dec
constexpr stream_manipulator dec
Definition: stream.hpp:679
cl::__ESIMD_ENS::detail::to_lsc_atomic_op
constexpr lsc_atomic_op to_lsc_atomic_op()
Definition: common.hpp:202
cl::__ESIMD_ENS::argument_type::U8
@ U8
cl::__ESIMD_ENS::lsc_fence_op::flushl3
@ flushl3
dirty lines are written to memory, but retained in cache in clean state
cl::__ESIMD_ENS::split_barrier_action::wait
@ wait
cl::__ESIMD_ENS::detail::lsc_atomic_op::iinc
@ iinc
cl::__ESIMD_ENS::argument_type::BF16
@ BF16
cl::__ESIMD_ENS::detail::lsc_vector_size::n16
@ n16
cl::__ESIMD_ENS::lsc_fence_op::evict
@ evict
no operation
cl::__ESIMD_ENS::detail::lsc_atomic_op::isub
@ isub
cl::__ESIMD_ENS::lsc_scope::tile
@ tile
flush out to the local scope
cl::__ESIMD_ENS::detail::lsc_vector_size
lsc_vector_size
Definition: common.hpp:112
cl::__ESIMD_ENS::detail::check_lsc_data_size
constexpr void check_lsc_data_size()
Definition: common.hpp:142
cl::__ESIMD_ENS::detail::lsc_expand_type::type
typename std::conditional< sizeof(T)< 4, uint32_t, T >::type type
Definition: common.hpp:322
cl::__ESIMD_ENS::detail::lsc_data_order::transpose
@ transpose
cl::__ESIMD_ENS::detail::finalize_data_size
constexpr lsc_data_size finalize_data_size()
Definition: common.hpp:297
cl::__ESIMD_ENS::cache_hint::streaming
@ streaming
cl::__ESIMD_ENS::cache_hint::uncached
@ uncached
cl::__ESIMD_ENS::detail::lsc_vector_size::n32
@ n32
cl::__ESIMD_ENS::lsc_data_size::u8u32
@ u8u32
cl::__ESIMD_ENS::detail::lsc_vector_size::n8
@ n8
cl::__ESIMD_ENS::argument_type::U1
@ U1
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12