DPC++ Runtime
Runtime libraries for oneAPI DPC++
operators.hpp
Go to the documentation of this file.
1 //==-------------- operators.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 // Binary operator definitions for ESIMD types.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
17 
20 
21 // Put operators into the ESIMD detail namespace to make argument-dependent
22 // lookup find these operators instead of those defined in e.g. sycl namespace
23 // (which would stop further lookup, leaving just non-viable sycl::operator <
24 // etc. on the table).
25 
27 namespace __ESIMD_DNS {
28 // clang-format off
31 
102 // clang-format on
103 
105 // simd_obj_impl global operators
107 
108 // ========= simd_obj_impl bitwise logic and arithmetic operators
109 
110 #define __ESIMD_DEF_SIMD_OBJ_IMPL_BIN_OP(BINOP, BINOP_ID, COND) \
111  \
112  /* simd_obj_impl BINOP simd_obj_impl */ \
113  template <class T1, class T2, int N, template <class, int> class SimdT, \
114  class SimdTx = SimdT<T1, N>, class = std::enable_if_t<COND>> \
115  inline auto operator BINOP( \
116  const __ESIMD_DNS::simd_obj_impl<__raw_t<T1>, N, SimdT<T1, N>> &LHS, \
117  const __ESIMD_DNS::simd_obj_impl<__raw_t<T2>, N, SimdT<T2, N>> &RHS) { \
118  if constexpr (__ESIMD_DNS::is_simd_type_v<SimdT<T1, N>>) { \
119  /* TODO CM_COMPAT To match CM behavior, here vector element types are */ \
120  /* promoted similar to C++ scalar type promotion, which is different */ \
121  /* from clang behavior when applying binary ops to vectors */ \
122  using PromotedT = __ESIMD_DNS::computation_type_t<T1, T2>; \
123  /* vector_binary_op returns SimdT<PromotedT, N>::raw_vector_type */ \
124  SimdT<PromotedT, N> Res = vector_binary_op<BINOP_ID, PromotedT, N>( \
125  __ESIMD_DNS::convert_vector<PromotedT, T1, N>(LHS.data()), \
126  __ESIMD_DNS::convert_vector<PromotedT, T2, N>(RHS.data())); \
127  return Res; \
128  } else { \
129  /* for SimdT=simd_mask_impl T1 and T2 are both equal to \
130  * simd_mask_elem_type */ \
131  return SimdT<T1, N>(LHS.data() BINOP RHS.data()); \
132  } \
133  } \
134  \
135  /* simd_obj_impl BINOP SCALAR */ \
136  template <class T1, int N1, template <class, int> class SimdT1, class T2, \
137  class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t<COND>> \
138  inline auto operator BINOP( \
139  const __ESIMD_DNS::simd_obj_impl<__raw_t<T1>, N1, SimdT1<T1, N1>> &LHS, \
140  T2 RHS) { \
141  if constexpr (__ESIMD_DNS::is_simd_type_v<SimdT1<T1, N1>>) { \
142  /* convert the SCALAR to vector type and reuse the basic operation over \
143  * simd objects */ \
144  return LHS BINOP SimdT1<T2, N1>(RHS); \
145  } else { \
146  /* SimdT1 is a mask, T1 is mask element type - convert RHS implicitly to \
147  * T1 */ \
148  return LHS BINOP SimdT1<T1, N1>(RHS); \
149  } \
150  } \
151  \
152  /* SCALAR BINOP simd_obj_impl */ \
153  template <class T1, class T2, int N2, template <class, int> class SimdT2, \
154  class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t<COND>> \
155  inline auto operator BINOP( \
156  T1 LHS, \
157  const __ESIMD_DNS::simd_obj_impl<__raw_t<T2>, N2, SimdT2<T2, N2>> \
158  &RHS) { \
159  if constexpr (__ESIMD_DNS::is_simd_type_v<SimdT2<T2, N2>>) { \
160  /* convert the SCALAR to vector type and reuse the basic operation over \
161  * simd objects */ \
162  return SimdT2<T1, N2>(LHS) BINOP RHS; \
163  } else { \
164  /* simd_mask_case */ \
165  return SimdT2<T2, N2>(LHS) BINOP RHS; \
166  } \
167  }
168 
169 // TODO add doxygen for individual overloads.
170 #define __ESIMD_BITWISE_OP_FILTER \
171  std::is_integral_v<T1> &&std::is_integral_v<T2>
175 #undef __ESIMD_BITWISE_OP_FILTER
176 
177 #define __ESIMD_SHIFT_OP_FILTER \
178  std::is_integral_v<T1> &&std::is_integral_v<T2> \
179  &&__ESIMD_DNS::is_simd_type_v<SimdTx>
183 #undef __ESIMD_SHIFT_OP_FILTER
184 
185 #define __ESIMD_ARITH_OP_FILTER \
186  __ESIMD_DNS::is_valid_simd_elem_type_v<T1> \
187  &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> \
188  &&__ESIMD_DNS::is_simd_type_v<SimdTx>
189 
194 #undef __ESIMD_ARITH_OP_FILTER
195 
196 #undef __ESIMD_DEF_SIMD_OBJ_IMPL_BIN_OP
197 
198 // ========= simd_obj_impl comparison operators
199 // Both simd and simd_mask will match simd_obj_impl argument when resolving
200 // operator overloads.
201 
202 #define __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(CMPOP, CMPOP_ID, COND) \
203  \
204  /* simd_obj_impl CMPOP simd_obj_impl */ \
205  template <class T1, class T2, int N, template <class, int> class SimdT, \
206  class SimdTx = SimdT<T1, N>, class = std::enable_if_t<COND>> \
207  inline __ESIMD_NS::simd_mask<N> operator CMPOP( \
208  const __ESIMD_DNS::simd_obj_impl<__raw_t<T1>, N, SimdT<T1, N>> &LHS, \
209  const __ESIMD_DNS::simd_obj_impl<__raw_t<T2>, N, SimdT<T2, N>> &RHS) { \
210  using MaskVecT = typename __ESIMD_NS::simd_mask<N>::raw_vector_type; \
211  \
212  if constexpr (is_simd_type_v<SimdT<T1, N>>) { \
213  using PromotedT = computation_type_t<T1, T2>; \
214  /* vector_comparison_op returns vector_type_t<Ti, N>, where Ti is \
215  * integer type */ \
216  /* of the same bit size as PromotedT */ \
217  auto Res = vector_comparison_op<CMPOP_ID, PromotedT, N>( \
218  __ESIMD_DNS::convert_vector<PromotedT, T1, N>(LHS.data()), \
219  __ESIMD_DNS::convert_vector<PromotedT, T2, N>(RHS.data())); \
220  using ResElemT = element_type_t<decltype(Res)>; \
221  return __ESIMD_NS::simd_mask<N>( \
222  __ESIMD_DNS::convert_vector<simd_mask_elem_type, ResElemT, N>(Res) & \
223  MaskVecT(1)); \
224  } else { \
225  /* this is comparison of masks, don't perform type promotion */ \
226  auto ResVec = LHS.data() CMPOP RHS.data(); \
227  return __ESIMD_NS::simd_mask<N>(__ESIMD_DNS::convert<MaskVecT>(ResVec) & \
228  MaskVecT(1)); \
229  } \
230  } \
231  \
232  /* simd_obj_impl CMPOP SCALAR */ \
233  template <class T1, int N1, template <class, int> class SimdT1, class T2, \
234  class SimdTx = SimdT1<T1, N1>, \
235  class = std::enable_if_t< \
236  __ESIMD_DNS::is_valid_simd_elem_type_v<T2> && COND>> \
237  inline __ESIMD_NS::simd_mask<N1> operator CMPOP( \
238  const __ESIMD_DNS::simd_obj_impl<__raw_t<T1>, N1, SimdT1<T1, N1>> &LHS, \
239  T2 RHS) { \
240  if constexpr (__ESIMD_DNS::is_simd_type_v<SimdT1<T1, N1>>) \
241  /* simd case */ \
242  return LHS CMPOP SimdT1<T2, N1>(RHS); \
243  else \
244  /* simd_mask case - element type is fixed */ \
245  return LHS CMPOP SimdT1<T1, N1>(convert_scalar<T1>(RHS)); \
246  } \
247  \
248  /* SCALAR CMPOP simd_obj_impl */ \
249  template <class T1, class T2, int N2, template <class, int> class SimdT2, \
250  class SimdTx = SimdT2<T2, N2>, \
251  class = std::enable_if_t< \
252  __ESIMD_DNS::is_valid_simd_elem_type_v<T1> && COND>> \
253  inline __ESIMD_NS::simd_mask<N2> operator CMPOP( \
254  T1 LHS, \
255  const __ESIMD_DNS::simd_obj_impl<__raw_t<T2>, N2, SimdT2<T2, N2>> \
256  &RHS) { \
257  if constexpr (__ESIMD_DNS::is_simd_type_v<SimdT2<T2, N2>>) \
258  /* simd case */ \
259  return SimdT2<T1, N2>(LHS) CMPOP RHS; \
260  else \
261  /* simd_mask case - element type is fixed */ \
262  return SimdT2<T2, N2>(convert_scalar<T2>(LHS)) CMPOP RHS; \
263  }
264 
265 // Equality comparison is defined for all simd_obj_impl subclasses.
266 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(==, CmpOp::eq, true)
267 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(!=, CmpOp::ne, true)
268 
269 // Relational operators are defined only for the simd type.
271  __ESIMD_DNS::is_simd_type_v<SimdTx>)
273  __ESIMD_DNS::is_simd_type_v<SimdTx>)
274 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(<=, CmpOp::lte,
275  __ESIMD_DNS::is_simd_type_v<SimdTx>)
276 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(>=, CmpOp::gte,
277  __ESIMD_DNS::is_simd_type_v<SimdTx>)
278 
279 // Logical operators are defined only for the simd_mask type
280 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(&&, BinOp::log_and,
281  __ESIMD_DNS::is_simd_mask_type_v<SimdTx>)
282 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(||, BinOp::log_or,
283  __ESIMD_DNS::is_simd_mask_type_v<SimdTx>)
284 
285 #undef __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP
286 
288 // simd_view global operators
290 
291 // ========= simd_view bitwise logic and arithmetic operators
292 
293 #define __ESIMD_DEF_SIMD_VIEW_BIN_OP(BINOP, COND) \
294  \
295  /* simd_view BINOP simd_view */ \
296  template < \
297  class SimdT1, class RegionT1, class SimdT2, class RegionT2, \
298  class T1 = typename __ESIMD_NS::shape_type<RegionT1>::element_type, \
299  class T2 = typename __ESIMD_NS::shape_type<RegionT2>::element_type, \
300  auto N1 = __ESIMD_NS::shape_type<RegionT1>::length, \
301  auto N2 = __ESIMD_NS::shape_type<RegionT2>::length, \
302  class = std::enable_if_t<__ESIMD_DNS::is_simd_type_v<SimdT1> == \
303  __ESIMD_DNS::is_simd_type_v<SimdT2> && \
304  (N1 == N2 || N1 == 1 || N2 == 1) && COND>> \
305  inline auto operator BINOP( \
306  const __ESIMD_NS::simd_view<SimdT1, RegionT1> &LHS, \
307  const __ESIMD_NS::simd_view<SimdT2, RegionT2> &RHS) { \
308  if constexpr (N1 == 1) \
309  return (T1)LHS.read()[0] BINOP RHS.read(); \
310  else if constexpr (N2 == 1) \
311  return LHS.read() BINOP(T2) RHS.read()[0]; \
312  else \
313  return LHS.read() BINOP RHS.read(); \
314  } \
315  \
316  /* simd* BINOP simd_view<simd*...> */ \
317  template <class SimdT1, class SimdT2, class RegionT2, \
318  class T1 = typename SimdT1::element_type, \
319  class T2 = \
320  typename __ESIMD_NS::shape_type<RegionT2>::element_type, \
321  class = std::enable_if_t< \
322  __ESIMD_DNS::is_simd_obj_impl_derivative_v<SimdT1> && \
323  (__ESIMD_DNS::is_simd_type_v<SimdT1> == \
324  __ESIMD_DNS::is_simd_type_v< \
325  SimdT2>)&&(SimdT1::length == \
326  __ESIMD_NS::shape_type<RegionT2>::length) && \
327  COND>> \
328  inline auto operator BINOP( \
329  const SimdT1 &LHS, const __ESIMD_NS::simd_view<SimdT2, RegionT2> &RHS) { \
330  return LHS BINOP RHS.read(); \
331  } \
332  \
333  /* simd_view<simd*...> BINOP simd* */ \
334  template < \
335  class SimdT1, class RegionT1, class SimdT2, \
336  class T1 = typename __ESIMD_NS::shape_type<RegionT1>::element_type, \
337  class T2 = typename SimdT2::element_type, \
338  class = std::enable_if_t< \
339  __ESIMD_DNS::is_simd_obj_impl_derivative_v<SimdT2> && \
340  __ESIMD_DNS::is_simd_type_v<SimdT1> == \
341  __ESIMD_DNS::is_simd_type_v<SimdT2> && \
342  (SimdT2::length == __ESIMD_NS::shape_type<RegionT1>::length) && \
343  COND>> \
344  inline auto operator BINOP( \
345  const __ESIMD_NS::simd_view<SimdT1, RegionT1> &LHS, const SimdT2 &RHS) { \
346  return LHS.read() BINOP RHS; \
347  } \
348  \
349  /* SCALAR BINOP simd_view */ \
350  template <class T1, class SimdViewT2, \
351  class T2 = typename SimdViewT2::element_type, \
352  class SimdT1 = typename SimdViewT2::value_type, \
353  class = std::enable_if_t< \
354  __ESIMD_DNS::is_any_simd_view_type_v<SimdViewT2> && COND>> \
355  inline auto operator BINOP(T1 LHS, const SimdViewT2 &RHS) { \
356  return LHS BINOP RHS.read(); \
357  } \
358  \
359  /* simd_view BINOP SCALAR */ \
360  template <class SimdViewT1, class T2, \
361  class T1 = typename SimdViewT1::element_type, \
362  class SimdT1 = typename SimdViewT1::value_type, \
363  class = std::enable_if_t< \
364  __ESIMD_DNS::is_any_simd_view_type_v<SimdViewT1> && COND>> \
365  inline auto operator BINOP(const SimdViewT1 &LHS, T2 RHS) { \
366  return LHS.read() BINOP RHS; \
367  }
368 
369 #define __ESIMD_BITWISE_OP_FILTER \
370  std::is_integral_v<T1> &&std::is_integral_v<T2>
374 #undef __ESIMD_BITWISE_OP_FILITER
375 
376 #define __ESIMD_SHIFT_OP_FILTER \
377  std::is_integral_v<T1> &&std::is_integral_v<T2> \
378  &&__ESIMD_DNS::is_simd_type_v<SimdT1>
379 
383 #undef __ESIMD_SHIFT_OP_FILTER
384 
385 #define __ESIMD_ARITH_OP_FILTER \
386  __ESIMD_DNS::is_simd_type_v<SimdT1> \
387  &&__ESIMD_DNS::is_valid_simd_elem_type_v<T1> \
388  &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2>
389 
394 #undef __ESIMD_ARITH_OP_FILTER
395 
396 __ESIMD_DEF_SIMD_VIEW_BIN_OP(&&, __ESIMD_DNS::is_simd_mask_type_v<SimdT1>)
397 __ESIMD_DEF_SIMD_VIEW_BIN_OP(||, __ESIMD_DNS::is_simd_mask_type_v<SimdT1>)
398 
399 #undef __ESIMD_DEF_SIMD_VIEW_BIN_OP
400 
401 // ========= simd_view comparison operators
402 
403 #define __ESIMD_DEF_SIMD_VIEW_CMP_OP(CMPOP, COND) \
404  \
405  /* simd_view CMPOP simd_view */ \
406  template <class SimdT1, class RegionT1, class SimdT2, class RegionT2, \
407  auto N1 = __ESIMD_NS::shape_type<RegionT1>::length, \
408  auto N2 = __ESIMD_NS::shape_type<RegionT2>::length, \
409  class = std::enable_if_t</* both views must have the same base \
410  type kind - simds or masks: */ \
411  (__ESIMD_DNS::is_simd_type_v<SimdT1> == \
412  __ESIMD_DNS::is_simd_type_v< \
413  SimdT2>)&&/* the length of the views \
414  must match as well: */ \
415  (N1 == N2 || N1 == 1 || N2 == 1) && \
416  COND>> \
417  inline auto operator CMPOP( \
418  const __ESIMD_NS::simd_view<SimdT1, RegionT1> &LHS, \
419  const __ESIMD_NS::simd_view<SimdT2, RegionT2> &RHS) { \
420  using T1 = typename __ESIMD_NS::shape_type<RegionT1>::element_type; \
421  using T2 = typename __ESIMD_NS::shape_type<RegionT2>::element_type; \
422  if constexpr (N1 == 1) \
423  return (T1)LHS.read()[0] CMPOP RHS.read(); \
424  else if constexpr (N2 == 1) \
425  return LHS.read() CMPOP(T2) RHS.read()[0]; \
426  else \
427  return LHS.read() CMPOP RHS.read(); \
428  } \
429  \
430  /* simd_view CMPOP simd_obj_impl */ \
431  template <class SimdT1, class RegionT1, class RawT2, int N2, class SimdT2, \
432  class = std::enable_if_t< \
433  (__ESIMD_NS::shape_type<RegionT1>::length == N2) && \
434  (__ESIMD_DNS::is_simd_type_v<SimdT1> == \
435  __ESIMD_DNS::is_simd_type_v<SimdT2>)&&COND>> \
436  inline __ESIMD_NS::simd_mask<N2> operator CMPOP( \
437  const __ESIMD_NS::simd_view<SimdT1, RegionT1> &LHS, \
438  const __ESIMD_DNS::simd_obj_impl<RawT2, N2, SimdT2> &RHS) { \
439  return LHS.read() CMPOP SimdT2(RHS.data()); \
440  } \
441  \
442  /* simd_obj_impl CMPOP simd_view */ \
443  template <class RawT1, int N1, class SimdT1, class SimdT2, class RegionT2, \
444  class = std::enable_if_t< \
445  (__ESIMD_NS::shape_type<RegionT2>::length == N1) && \
446  (__ESIMD_DNS::is_simd_type_v<SimdT1> == \
447  __ESIMD_DNS::is_simd_type_v<SimdT2>)&&COND>> \
448  inline __ESIMD_NS::simd_mask<N1> operator CMPOP( \
449  const __ESIMD_DNS::simd_obj_impl<RawT1, N1, SimdT1> &LHS, \
450  const __ESIMD_NS::simd_view<SimdT2, RegionT2> &RHS) { \
451  return SimdT1(LHS.data()) CMPOP RHS.read(); \
452  } \
453  \
454  /* simd_view CMPOP SCALAR */ \
455  template <class SimdT1, class RegionT1, class T2, \
456  class = std::enable_if_t< \
457  __ESIMD_DNS::is_valid_simd_elem_type_v<T2> && COND>> \
458  inline auto operator CMPOP( \
459  const __ESIMD_NS::simd_view<SimdT1, RegionT1> &LHS, T2 RHS) { \
460  return LHS.read() CMPOP RHS; \
461  } \
462  \
463  /* SCALAR CMPOP simd_view */ \
464  template <class T1, class SimdT2, class RegionT2, class SimdT1 = SimdT2, \
465  class = std::enable_if_t< \
466  __ESIMD_DNS::is_valid_simd_elem_type_v<T1> && COND>> \
467  inline auto operator CMPOP( \
468  T1 LHS, const __ESIMD_NS::simd_view<SimdT2, RegionT2> &RHS) { \
469  return LHS CMPOP RHS.read(); \
470  }
471 
472 // Equality comparison is defined for views of all simd_obj_impl derivatives.
475 
476 // Relational operators are defined only for views of the simd class.
477 __ESIMD_DEF_SIMD_VIEW_CMP_OP(<, __ESIMD_DNS::is_simd_type_v<SimdT1>)
478 __ESIMD_DEF_SIMD_VIEW_CMP_OP(>, __ESIMD_DNS::is_simd_type_v<SimdT1>)
479 __ESIMD_DEF_SIMD_VIEW_CMP_OP(<=, __ESIMD_DNS::is_simd_type_v<SimdT1>)
480 __ESIMD_DEF_SIMD_VIEW_CMP_OP(>=, __ESIMD_DNS::is_simd_type_v<SimdT1>)
481 
482 #undef __ESIMD_DEF_SIMD_VIEW_CMP_OP
483 
484 } // namespace __ESIMD_DNS
485 } // __SYCL_INLINE_NAMESPACE(cl)
__ESIMD_DEF_SIMD_VIEW_CMP_OP
#define __ESIMD_DEF_SIMD_VIEW_CMP_OP(CMPOP, COND)
cl::sycl::bit_xor
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
cl::sycl::detail::device_global_map::add
void add(const void *DeviceGlobalPtr, const char *UniqueId)
Definition: device_global_map.cpp:16
simd_view.hpp
__ESIMD_DEF_SIMD_VIEW_BIN_OP
#define __ESIMD_DEF_SIMD_VIEW_BIN_OP(BINOP, COND)
cl::sycl::bit_or
std::bit_or< T > bit_or
Definition: functional.hpp:21
__ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP
#define __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(CMPOP, CMPOP_ID, COND)
Definition: operators.hpp:198
__ESIMD_SHIFT_OP_FILTER
#define __ESIMD_SHIFT_OP_FILTER
Definition: operators.hpp:173
cl::sycl::ext::intel::experimental::esimd::shl
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > > shl(sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
Shift left operation (vector version)
Definition: math.hpp:40
simd_obj_impl.hpp
__ESIMD_BITWISE_OP_FILTER
#define __ESIMD_BITWISE_OP_FILTER
Definition: operators.hpp:166
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
simd.hpp
elem_type_traits.hpp
types.hpp
simd_view_impl.hpp
__ESIMD_DEF_SIMD_OBJ_IMPL_BIN_OP
#define __ESIMD_DEF_SIMD_OBJ_IMPL_BIN_OP(BINOP, BINOP_ID, COND)
Definition: operators.hpp:110
cl::sycl::ext::intel::experimental::esimd::shr
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > > shr(sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
Shift right operation (vector version)
Definition: math.hpp:114
cl::sycl::ext::intel::experimental::esimd::div
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T, SZ > > div(sycl::ext::intel::esimd::simd< T, SZ > &remainder, sycl::ext::intel::esimd::simd< T, SZ > src0, U src1)
Integral division with a vector dividend and a scalar divisor.
Definition: math.hpp:541
__ESIMD_ARITH_OP_FILTER
#define __ESIMD_ARITH_OP_FILTER
Definition: operators.hpp:181
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12