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