DPC++ Runtime
Runtime libraries for oneAPI DPC++
vector.hpp
Go to the documentation of this file.
1 //==---------------- vector.hpp --- Implements sycl::vec -------------------==//
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 
9 #pragma once
10 
11 // Check if Clang's ext_vector_type attribute is available. Host compiler
12 // may not be Clang, and Clang may not be built with the extension.
13 #ifdef __clang__
14 #ifndef __has_extension
15 #define __has_extension(x) 0
16 #endif
17 #ifdef __HAS_EXT_VECTOR_TYPE__
18 #error "Undefine __HAS_EXT_VECTOR_TYPE__ macro"
19 #endif
20 #if __has_extension(attribute_ext_vector_type)
21 #define __HAS_EXT_VECTOR_TYPE__
22 #endif
23 #endif // __clang__
24 
25 #if !defined(__HAS_EXT_VECTOR_TYPE__) && defined(__SYCL_DEVICE_ONLY__)
26 #error "SYCL device compiler is built without ext_vector_type support"
27 #endif
28 
29 #include <sycl/access/access.hpp> // for decorated, address_space
30 #include <sycl/aliases.hpp> // for half, cl_char, cl_int
31 #include <sycl/detail/common.hpp> // for ArrayCreator, RepeatV...
32 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
33 #include <sycl/detail/generic_type_lists.hpp> // for vector_basic_list
34 #include <sycl/detail/generic_type_traits.hpp> // for is_sigeninteger, is_s...
35 #include <sycl/detail/memcpy.hpp> // for memcpy
36 #include <sycl/detail/type_list.hpp> // for is_contained
37 #include <sycl/detail/type_traits.hpp> // for is_floating_point
38 #include <sycl/detail/vector_arith.hpp>
39 #include <sycl/half_type.hpp> // for StorageT, half, Vec16...
40 
41 #include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16
42 
43 #include <algorithm> // for std::min
44 #include <array> // for array
45 #include <cassert> // for assert
46 #include <cstddef> // for size_t, NULL, byte
47 #include <cstdint> // for uint8_t, int16_t, int...
48 #include <functional> // for divides, multiplies
49 #include <iterator> // for pair
50 #include <ostream> // for operator<<, basic_ost...
51 #include <type_traits> // for enable_if_t, is_same
52 #include <utility> // for index_sequence, make_...
53 
54 namespace sycl {
55 
56 // TODO: Fix in the next ABI breaking windows.
57 enum class rounding_mode { automatic = 0, rte = 1, rtz = 2, rtp = 3, rtn = 4 };
58 
59 inline namespace _V1 {
60 
61 struct elem {
62  static constexpr int x = 0;
63  static constexpr int y = 1;
64  static constexpr int z = 2;
65  static constexpr int w = 3;
66  static constexpr int r = 0;
67  static constexpr int g = 1;
68  static constexpr int b = 2;
69  static constexpr int a = 3;
70  static constexpr int s0 = 0;
71  static constexpr int s1 = 1;
72  static constexpr int s2 = 2;
73  static constexpr int s3 = 3;
74  static constexpr int s4 = 4;
75  static constexpr int s5 = 5;
76  static constexpr int s6 = 6;
77  static constexpr int s7 = 7;
78  static constexpr int s8 = 8;
79  static constexpr int s9 = 9;
80  static constexpr int sA = 10;
81  static constexpr int sB = 11;
82  static constexpr int sC = 12;
83  static constexpr int sD = 13;
84  static constexpr int sE = 14;
85  static constexpr int sF = 15;
86 };
87 
88 namespace detail {
89 template <typename VecT, typename OperationLeftT, typename OperationRightT,
90  template <typename> class OperationCurrentT, int... Indexes>
91 class SwizzleOp;
92 
93 // Special type indicating that SwizzleOp should just read value from vector -
94 // not trying to perform any operations. Should not be called.
95 template <typename T> class GetOp {
96 public:
97  using DataT = T;
98  DataT getValue(size_t) const { return (DataT)0; }
99  DataT operator()(DataT, DataT) { return (DataT)0; }
100 };
101 
102 // Templated vs. non-templated conversion operator behaves differently when two
103 // conversions are needed as in the case below:
104 //
105 // sycl::vec<int, 1> v;
106 // std::ignore = static_cast<bool>(v);
107 //
108 // Make sure the snippet above compiles. That is important because
109 //
110 // sycl::vec<int, 2> v;
111 // if (v.x() == 42)
112 // ...
113 //
114 // must go throw `v.x()` returning a swizzle, then its `operator==` returning
115 // vec<int, 1> and we want that code to compile.
116 template <typename Vec, typename T, int N, typename = void>
117 struct ScalarConversionOperatorMixIn {};
118 
119 template <typename Vec, typename T, int N>
120 struct ScalarConversionOperatorMixIn<Vec, T, N, std::enable_if_t<N == 1>> {
121  operator T() const { return (*static_cast<const Vec *>(this))[0]; }
122 };
123 
124 } // namespace detail
125 
127 // Provides a cross-platform vector class template that works efficiently on
128 // SYCL devices as well as in host C++ code.
129 template <typename DataT, int NumElements>
130 class __SYCL_EBO vec
131  : public detail::vec_arith<DataT, NumElements>,
132  public detail::ScalarConversionOperatorMixIn<vec<DataT, NumElements>,
133  DataT, NumElements> {
134 
135  static_assert(NumElements == 1 || NumElements == 2 || NumElements == 3 ||
136  NumElements == 4 || NumElements == 8 || NumElements == 16,
137  "Invalid number of elements for sycl::vec: only 1, 2, 3, 4, 8 "
138  "or 16 are supported");
139  static_assert(sizeof(bool) == sizeof(uint8_t), "bool size is not 1 byte");
140 
141  // https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#memory-layout-and-alignment
142  // It is required by the SPEC to align vec<DataT, 3> with vec<DataT, 4>.
143  static constexpr size_t AdjustedNum = (NumElements == 3) ? 4 : NumElements;
144 
145  // This represent type of underlying value. There should be only one field
146  // in the class, so vec<float, 16> should be equal to float16 in memory.
147  using DataType = std::array<DataT, AdjustedNum>;
148 
149 #ifdef __SYCL_DEVICE_ONLY__
150  using element_type_for_vector_t = typename detail::map_type<
151  DataT,
152 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
153  std::byte, /*->*/ std::uint8_t, //
154 #endif
155  bool, /*->*/ std::uint8_t, //
156  sycl::half, /*->*/ sycl::detail::half_impl::StorageT, //
157  sycl::ext::oneapi::bfloat16,
158  /*->*/ sycl::ext::oneapi::detail::Bfloat16StorageT, //
159  char, /*->*/ detail::ConvertToOpenCLType_t<char>, //
160  DataT, /*->*/ DataT //
161  >::type;
162 
163 public:
164  // Type used for passing sycl::vec to SPIRV builtins.
165  // We can not use ext_vector_type(1) as it's not supported by SPIRV
166  // plugins (CTS fails).
167  using vector_t =
168  typename std::conditional_t<NumElements == 1, element_type_for_vector_t,
169  element_type_for_vector_t __attribute__((
170  ext_vector_type(NumElements)))>;
171 
172 private:
173 #endif // __SYCL_DEVICE_ONLY__
174 
175  static constexpr int getNumElements() { return NumElements; }
176 
177  // SizeChecker is needed for vec(const argTN &... args) ctor to validate args.
178  template <int Counter, int MaxValue, class...>
179  struct SizeChecker : std::conditional_t<Counter == MaxValue, std::true_type,
180  std::false_type> {};
181 
182  template <int Counter, int MaxValue, typename DataT_, class... tail>
183  struct SizeChecker<Counter, MaxValue, DataT_, tail...>
184  : std::conditional_t<Counter + 1 <= MaxValue,
185  SizeChecker<Counter + 1, MaxValue, tail...>,
186  std::false_type> {};
187 
188  // Utility trait for creating an std::array from an vector argument.
189  template <typename DataT_, typename T> class FlattenVecArg {
190  template <std::size_t... Is>
191  static constexpr auto helper(const T &V, std::index_sequence<Is...>) {
192  // FIXME: Swizzle's `operator[]` for expression trees seems to be broken
193  // and returns values of the underlying vector of some of the operands. On
194  // the other hand, `getValue()` gives correct results. This can be changed
195  // to using `operator[]` once the bug is fixed.
196  if constexpr (detail::is_swizzle_v<T>)
197  return std::array{static_cast<DataT_>(V.getValue(Is))...};
198  else
199  return std::array{static_cast<DataT_>(V[Is])...};
200  }
201 
202  public:
203  constexpr auto operator()(const T &A) const {
204  if constexpr (detail::is_vec_or_swizzle_v<T>) {
205  return helper(A, std::make_index_sequence<T ::size()>());
206  } else {
207  return std::array{static_cast<DataT_>(A)};
208  }
209  }
210  };
211 
212  // Alias for shortening the vec arguments to array converter.
213  template <typename DataT_, typename... ArgTN>
214  using VecArgArrayCreator =
215  detail::ArrayCreator<DataT_, FlattenVecArg, ArgTN...>;
216 
217  template <int... Indexes>
218  using Swizzle =
219  detail::SwizzleOp<vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
220  detail::GetOp, Indexes...>;
221 
222  template <int... Indexes>
223  using ConstSwizzle =
224  detail::SwizzleOp<const vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
225  detail::GetOp, Indexes...>;
226 
227  // Shortcuts for args validation in vec(const argTN &... args) ctor.
228  template <typename CtorArgTy>
229  static constexpr bool AllowArgTypeInVariadicCtor = []() constexpr {
230  // FIXME: This logic implements the behavior of the previous implementation.
231  if constexpr (detail::is_vec_or_swizzle_v<CtorArgTy>) {
232  if constexpr (CtorArgTy::size() == 1)
233  return std::is_convertible_v<typename CtorArgTy::element_type, DataT>;
234  else
235  return std::is_same_v<typename CtorArgTy::element_type, DataT>;
236  } else {
237  return std::is_convertible_v<CtorArgTy, DataT>;
238  }
239  }();
240 
241  template <typename T> static constexpr int num_elements() {
242  if constexpr (detail::is_vec_or_swizzle_v<T>)
243  return T::size();
244  else
245  return 1;
246  }
247 
248  // Element type for relational operator return value.
249  using rel_t = detail::select_cl_scalar_integral_signed_t<DataT>;
250 
251 public:
252  // Aliases required by SYCL 2020 to make sycl::vec consistent
253  // with that of marray and buffer.
254  using element_type = DataT;
255  using value_type = DataT;
256 
257  /****************** Constructors **************/
258  vec() = default;
259  constexpr vec(const vec &Rhs) = default;
260  constexpr vec(vec &&Rhs) = default;
261 
262 private:
263  // Implementation detail for the next public ctor.
264  template <size_t... Is>
265  constexpr vec(const std::array<DataT, NumElements> &Arr,
266  std::index_sequence<Is...>)
267  : m_Data{Arr[Is]...} {}
268 
269 public:
270  explicit constexpr vec(const DataT &arg)
271  : vec{detail::RepeatValue<NumElements>(arg),
272  std::make_index_sequence<NumElements>()} {}
273 
274  // Constructor from values of base type or vec of base type. Checks that
275  // base types are match and that the NumElements == sum of lengths of args.
276  template <typename... argTN,
277  typename = std::enable_if_t<
278  ((AllowArgTypeInVariadicCtor<argTN> && ...)) &&
279  ((num_elements<argTN>() + ...)) == NumElements>>
280  constexpr vec(const argTN &...args)
281  : vec{VecArgArrayCreator<DataT, argTN...>::Create(args...),
282  std::make_index_sequence<NumElements>()} {}
283 
284  /****************** Assignment Operators **************/
285  constexpr vec &operator=(const vec &Rhs) = default;
286 
287  // Template required to prevent ambiguous overload with the copy assignment
288  // when NumElements == 1. The template prevents implicit conversion from
289  // vec<_, 1> to DataT.
290  template <typename Ty = DataT>
291  typename std::enable_if_t<
292  std::is_fundamental_v<Ty> ||
293  detail::is_half_or_bf16_v<typename std::remove_const_t<Ty>>,
294  vec &>
295  operator=(const DataT &Rhs) {
296  *this = vec{Rhs};
297  return *this;
298  }
299 
300  // W/o this, things like "vec<char,*> = vec<signed char, *>" doesn't work.
301  template <typename Ty = DataT>
302  typename std::enable_if_t<
303  !std::is_same_v<Ty, rel_t> && std::is_convertible_v<Ty, rel_t>, vec &>
304  operator=(const vec<rel_t, NumElements> &Rhs) {
305  *this = Rhs.template as<vec>();
306  return *this;
307  }
308 
309 #ifdef __SYCL_DEVICE_ONLY__
310  // Make it a template to avoid ambiguity with `vec(const DataT &)` when
311  // `vector_t` is the same as `DataT`. Not that the other ctor isn't a template
312  // so we don't even need a smart `enable_if` condition here, the mere fact of
313  // this being a template makes the other ctor preferred.
314  template <
315  typename vector_t_ = vector_t,
316  typename = typename std::enable_if_t<std::is_same_v<vector_t_, vector_t>>>
317  constexpr vec(vector_t_ openclVector) {
318  m_Data = sycl::bit_cast<DataType>(openclVector);
319  }
320 
321  /* @SYCL2020
322  * Available only when: compiled for the device.
323  * Converts this SYCL vec instance to the underlying backend-native vector
324  * type defined by vector_t.
325  */
326  operator vector_t() const { return sycl::bit_cast<vector_t>(m_Data); }
327 #endif // __SYCL_DEVICE_ONLY__
328 
329  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
330  static constexpr size_t get_count() { return size(); }
331  static constexpr size_t size() noexcept { return NumElements; }
332  __SYCL2020_DEPRECATED(
333  "get_size() is deprecated, please use byte_size() instead")
334  static constexpr size_t get_size() { return byte_size(); }
335  static constexpr size_t byte_size() noexcept { return sizeof(m_Data); }
336 
337 private:
338  // getValue should be able to operate on different underlying
339  // types: enum cl_float#N , builtin vector float#N, builtin type float.
340  constexpr auto getValue(int Index) const {
341  using RetType =
342  typename std::conditional_t<detail::is_byte_v<DataT>, int8_t,
343 #ifdef __SYCL_DEVICE_ONLY__
344  element_type_for_vector_t
345 #else
346  DataT
347 #endif
348  >;
349 
350 #ifdef __SYCL_DEVICE_ONLY__
351  if constexpr (std::is_same_v<DataT, sycl::ext::oneapi::bfloat16>)
352  return sycl::bit_cast<RetType>(m_Data[Index]);
353  else
354 #endif
355  return static_cast<RetType>(m_Data[Index]);
356  }
357 
358 public:
359  // Out-of-class definition is in `sycl/detail/vector_convert.hpp`
360  template <typename convertT,
361  rounding_mode roundingMode = rounding_mode::automatic>
362  vec<convertT, NumElements> convert() const;
363 
364  template <typename asT> asT as() const { return sycl::bit_cast<asT>(*this); }
365 
366  template <int... SwizzleIndexes> Swizzle<SwizzleIndexes...> swizzle() {
367  return this;
368  }
369 
370  template <int... SwizzleIndexes>
371  ConstSwizzle<SwizzleIndexes...> swizzle() const {
372  return this;
373  }
374 
375  const DataT &operator[](int i) const { return m_Data[i]; }
376 
377  DataT &operator[](int i) { return m_Data[i]; }
378 
379  // Begin hi/lo, even/odd, xyzw, and rgba swizzles. @{
380 private:
381  // Indexer used in the swizzles.def
382  // Currently it is defined as a template struct. Replacing it with a constexpr
383  // function would activate a bug in MSVC that is fixed only in v19.20.
384  // Until then MSVC does not recognize such constexpr functions as const and
385  // thus does not let using them in template parameters inside swizzle.def.
386  template <int Index> struct Indexer {
387  static constexpr int value = Index;
388  };
389 
390 public:
391 #ifdef __SYCL_ACCESS_RETURN
392 #error "Undefine __SYCL_ACCESS_RETURN macro"
393 #endif
394 #define __SYCL_ACCESS_RETURN this
395 #include "swizzles.def"
396 #undef __SYCL_ACCESS_RETURN
397  // }@ End of hi/lo, even/odd, xyzw, and rgba swizzles.
398 
399  template <access::address_space Space, access::decorated DecorateAddress>
400  void load(size_t Offset, multi_ptr<const DataT, Space, DecorateAddress> Ptr) {
401  for (int I = 0; I < NumElements; I++) {
402  m_Data[I] = *multi_ptr<const DataT, Space, DecorateAddress>(
403  Ptr + Offset * NumElements + I);
404  }
405  }
406  template <access::address_space Space, access::decorated DecorateAddress>
407  void load(size_t Offset, multi_ptr<DataT, Space, DecorateAddress> Ptr) {
408  multi_ptr<const DataT, Space, DecorateAddress> ConstPtr(Ptr);
409  load(Offset, ConstPtr);
410  }
411  template <int Dimensions, access::mode Mode,
412  access::placeholder IsPlaceholder, access::target Target,
413  typename PropertyListT>
414  void
415  load(size_t Offset,
416  accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
417  Acc) {
418  multi_ptr<const DataT, detail::TargetToAS<Target>::AS,
419  access::decorated::yes>
420  MultiPtr(Acc);
421  load(Offset, MultiPtr);
422  }
423  void load(size_t Offset, const DataT *Ptr) {
424  for (int I = 0; I < NumElements; ++I)
425  m_Data[I] = Ptr[Offset * NumElements + I];
426  }
427 
428  template <access::address_space Space, access::decorated DecorateAddress>
429  void store(size_t Offset,
430  multi_ptr<DataT, Space, DecorateAddress> Ptr) const {
431  for (int I = 0; I < NumElements; I++) {
432  *multi_ptr<DataT, Space, DecorateAddress>(Ptr + Offset * NumElements +
433  I) = m_Data[I];
434  }
435  }
436  template <int Dimensions, access::mode Mode,
437  access::placeholder IsPlaceholder, access::target Target,
438  typename PropertyListT>
439  void
440  store(size_t Offset,
441  accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
442  Acc) {
443  multi_ptr<DataT, detail::TargetToAS<Target>::AS, access::decorated::yes>
444  MultiPtr(Acc);
445  store(Offset, MultiPtr);
446  }
447  void store(size_t Offset, DataT *Ptr) const {
448  for (int I = 0; I < NumElements; ++I)
449  Ptr[Offset * NumElements + I] = m_Data[I];
450  }
451 
452 private:
453  // fields
454  // Alignment is the same as size, to a maximum size of 64. SPEC requires
455  // "The elements of an instance of the SYCL vec class template are stored
456  // in memory sequentially and contiguously and are aligned to the size of
457  // the element type in bytes multiplied by the number of elements."
458  static constexpr int alignment = (std::min)((size_t)64, sizeof(DataType));
459  alignas(alignment) DataType m_Data;
460 
461  // friends
462  template <typename T1, typename T2, typename T3, template <typename> class T4,
463  int... T5>
464  friend class detail::SwizzleOp;
465  template <typename T1, int T2> friend class __SYCL_EBO vec;
466  // To allow arithmetic operators access private members of vec.
467  template <typename T1, int T2> friend class detail::vec_arith;
468  template <typename T1, int T2> friend class detail::vec_arith_common;
469 };
471 
472 #ifdef __cpp_deduction_guides
473 // all compilers supporting deduction guides also support fold expressions
474 template <class T, class... U,
475  class = std::enable_if_t<(std::is_same_v<T, U> && ...)>>
476 vec(T, U...) -> vec<T, sizeof...(U) + 1>;
477 #endif
478 
479 namespace detail {
480 
481 // Special type for working SwizzleOp with scalars, stores a scalar and gives
482 // the scalar at any index. Provides interface is compatible with SwizzleOp
483 // operations
484 template <typename T> class GetScalarOp {
485 public:
486  using DataT = T;
487  GetScalarOp(DataT Data) : m_Data(Data) {}
488  DataT getValue(size_t) const { return m_Data; }
489 
490 private:
491  DataT m_Data;
492 };
493 template <typename T>
494 using rel_t = detail::select_cl_scalar_integral_signed_t<T>;
495 
496 template <typename T> struct EqualTo {
497  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
498  return (Lhs == Rhs) ? -1 : 0;
499  }
500 };
501 
502 template <typename T> struct NotEqualTo {
503  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
504  return (Lhs != Rhs) ? -1 : 0;
505  }
506 };
507 
508 template <typename T> struct GreaterEqualTo {
509  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
510  return (Lhs >= Rhs) ? -1 : 0;
511  }
512 };
513 
514 template <typename T> struct LessEqualTo {
515  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
516  return (Lhs <= Rhs) ? -1 : 0;
517  }
518 };
519 
520 template <typename T> struct GreaterThan {
521  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
522  return (Lhs > Rhs) ? -1 : 0;
523  }
524 };
525 
526 template <typename T> struct LessThan {
527  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
528  return (Lhs < Rhs) ? -1 : 0;
529  }
530 };
531 
532 template <typename T> struct LogicalAnd {
533  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
534  return (Lhs && Rhs) ? -1 : 0;
535  }
536 };
537 
538 template <typename T> struct LogicalOr {
539  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
540  return (Lhs || Rhs) ? -1 : 0;
541  }
542 };
543 
544 template <typename T> struct RShift {
545  constexpr T operator()(const T &Lhs, const T &Rhs) const {
546  return Lhs >> Rhs;
547  }
548 };
549 
550 template <typename T> struct LShift {
551  constexpr T operator()(const T &Lhs, const T &Rhs) const {
552  return Lhs << Rhs;
553  }
554 };
555 
557 // SwizzleOP represents expression templates that operate on vec.
558 // Actual computation performed on conversion or assignment operators.
559 template <typename VecT, typename OperationLeftT, typename OperationRightT,
560  template <typename> class OperationCurrentT, int... Indexes>
561 class SwizzleOp {
562  using DataT = typename VecT::element_type;
563  // Certain operators return a vector with a different element type. Also, the
564  // left and right operand types may differ. CommonDataT selects a result type
565  // based on these types to ensure that the result value can be represented.
566  //
567  // Example 1:
568  // sycl::vec<unsigned char, 4> vec{...};
569  // auto result = 300u + vec.x();
570  //
571  // CommonDataT is std::common_type_t<OperationLeftT, OperationRightT> since
572  // it's larger than unsigned char.
573  //
574  // Example 2:
575  // sycl::vec<bool, 1> vec{...};
576  // auto result = vec.template swizzle<sycl::elem::s0>() && vec;
577  //
578  // CommonDataT is DataT since operator&& returns a vector with element type
579  // int8_t, which is larger than bool.
580  //
581  // Example 3:
582  // sycl::vec<std::byte, 4> vec{...}; auto swlo = vec.lo();
583  // auto result = swlo == swlo;
584  //
585  // CommonDataT is DataT since operator== returns a vector with element type
586  // int8_t, which is the same size as std::byte. std::common_type_t<DataT, ...>
587  // can't be used here since there's no type that int8_t and std::byte can both
588  // be implicitly converted to.
589  using OpLeftDataT = typename OperationLeftT::DataT;
590  using OpRightDataT = typename OperationRightT::DataT;
591  using CommonDataT = std::conditional_t<
592  sizeof(DataT) >= sizeof(std::common_type_t<OpLeftDataT, OpRightDataT>),
593  DataT, std::common_type_t<OpLeftDataT, OpRightDataT>>;
594  static constexpr int getNumElements() { return sizeof...(Indexes); }
595 
596  using rel_t = detail::rel_t<DataT>;
597  using vec_t = vec<DataT, sizeof...(Indexes)>;
598  using vec_rel_t = vec<rel_t, sizeof...(Indexes)>;
599 
600  template <typename OperationRightT_,
601  template <typename> class OperationCurrentT_, int... Idx_>
602  using NewLHOp = SwizzleOp<VecT,
603  SwizzleOp<VecT, OperationLeftT, OperationRightT,
604  OperationCurrentT, Indexes...>,
605  OperationRightT_, OperationCurrentT_, Idx_...>;
606 
607  template <typename OperationRightT_,
608  template <typename> class OperationCurrentT_, int... Idx_>
609  using NewRelOp = SwizzleOp<vec<rel_t, VecT::getNumElements()>,
610  SwizzleOp<VecT, OperationLeftT, OperationRightT,
611  OperationCurrentT, Indexes...>,
612  OperationRightT_, OperationCurrentT_, Idx_...>;
613 
614  template <typename OperationLeftT_,
615  template <typename> class OperationCurrentT_, int... Idx_>
616  using NewRHOp = SwizzleOp<VecT, OperationLeftT_,
617  SwizzleOp<VecT, OperationLeftT, OperationRightT,
618  OperationCurrentT, Indexes...>,
619  OperationCurrentT_, Idx_...>;
620 
621  template <int IdxNum, typename T = void>
622  using EnableIfOneIndex = typename std::enable_if_t<
623  1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
624 
625  template <int IdxNum, typename T = void>
626  using EnableIfMultipleIndexes = typename std::enable_if_t<
627  1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
628 
629  template <typename T>
630  using EnableIfScalarType = typename std::enable_if_t<
631  std::is_convertible_v<DataT, T> &&
632  (std::is_fundamental_v<T> ||
633  detail::is_half_or_bf16_v<typename std::remove_const_t<T>>)>;
634 
635  template <typename T>
636  using EnableIfNoScalarType = typename std::enable_if_t<
637  !std::is_convertible_v<DataT, T> ||
638  !(std::is_fundamental_v<T> ||
639  detail::is_half_or_bf16_v<typename std::remove_const_t<T>>)>;
640 
641  template <int... Indices>
642  using Swizzle =
643  SwizzleOp<VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
644 
645  template <int... Indices>
646  using ConstSwizzle =
647  SwizzleOp<const VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
648 
649 public:
650  using element_type = DataT;
651  using value_type = DataT;
652 
653 #ifdef __SYCL_DEVICE_ONLY__
654  using vector_t = typename vec_t::vector_t;
655 #endif // __SYCL_DEVICE_ONLY__
656 
657  const DataT &operator[](int i) const {
658  std::array<int, getNumElements()> Idxs{Indexes...};
659  return (*m_Vector)[Idxs[i]];
660  }
661 
662  template <typename _T = VecT>
663  std::enable_if_t<!std::is_const_v<_T>, DataT> &operator[](int i) {
664  std::array<int, getNumElements()> Idxs{Indexes...};
665  return (*m_Vector)[Idxs[i]];
666  }
667 
668  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
669  size_t get_count() const { return size(); }
670  static constexpr size_t size() noexcept { return getNumElements(); }
671 
672  template <int Num = getNumElements()>
673  __SYCL2020_DEPRECATED(
674  "get_size() is deprecated, please use byte_size() instead")
675  size_t get_size() const {
676  return byte_size<Num>();
677  }
678 
679  template <int Num = getNumElements()> size_t byte_size() const noexcept {
680  return sizeof(DataT) * (Num == 3 ? 4 : Num);
681  }
682 
683  template <typename T, int IdxNum = getNumElements(),
684  typename = EnableIfOneIndex<IdxNum>,
685  typename = EnableIfScalarType<T>>
686  operator T() const {
687  return getValue(0);
688  }
689 
690  template <typename T, typename = EnableIfScalarType<T>>
691  friend NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>
692  operator*(const T &Lhs, const SwizzleOp &Rhs) {
693  return NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
694  Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
695  }
696 
697  template <typename T, typename = EnableIfScalarType<T>>
698  friend NewRHOp<GetScalarOp<T>, std::plus, Indexes...>
699  operator+(const T &Lhs, const SwizzleOp &Rhs) {
700  return NewRHOp<GetScalarOp<T>, std::plus, Indexes...>(
701  Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
702  }
703 
704  template <typename T, typename = EnableIfScalarType<T>>
705  friend NewRHOp<GetScalarOp<T>, std::divides, Indexes...>
706  operator/(const T &Lhs, const SwizzleOp &Rhs) {
707  return NewRHOp<GetScalarOp<T>, std::divides, Indexes...>(
708  Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
709  }
710 
711  // TODO: Check that Rhs arg is suitable.
712 #ifdef __SYCL_OPASSIGN
713 #error "Undefine __SYCL_OPASSIGN macro."
714 #endif
715 #define __SYCL_OPASSIGN(OPASSIGN, OP) \
716  friend const SwizzleOp &operator OPASSIGN(const SwizzleOp & Lhs, \
717  const DataT & Rhs) { \
718  Lhs.operatorHelper<OP>(vec_t(Rhs)); \
719  return Lhs; \
720  } \
721  template <typename RhsOperation> \
722  friend const SwizzleOp &operator OPASSIGN(const SwizzleOp & Lhs, \
723  const RhsOperation & Rhs) { \
724  Lhs.operatorHelper<OP>(Rhs); \
725  return Lhs; \
726  } \
727  friend const SwizzleOp &operator OPASSIGN(const SwizzleOp & Lhs, \
728  const vec_t & Rhs) { \
729  Lhs.operatorHelper<OP>(Rhs); \
730  return Lhs; \
731  }
732 
733  __SYCL_OPASSIGN(+=, std::plus)
734  __SYCL_OPASSIGN(-=, std::minus)
735  __SYCL_OPASSIGN(*=, std::multiplies)
736  __SYCL_OPASSIGN(/=, std::divides)
737  __SYCL_OPASSIGN(%=, std::modulus)
738  __SYCL_OPASSIGN(&=, std::bit_and)
739  __SYCL_OPASSIGN(|=, std::bit_or)
740  __SYCL_OPASSIGN(^=, std::bit_xor)
741  __SYCL_OPASSIGN(>>=, RShift)
742  __SYCL_OPASSIGN(<<=, LShift)
743 #undef __SYCL_OPASSIGN
744 
745 #ifdef __SYCL_UOP
746 #error "Undefine __SYCL_UOP macro"
747 #endif
748 #define __SYCL_UOP(UOP, OPASSIGN) \
749  friend const SwizzleOp &operator UOP(const SwizzleOp & sv) { \
750  sv OPASSIGN static_cast<DataT>(1); \
751  return sv; \
752  } \
753  friend vec_t operator UOP(const SwizzleOp &sv, int) { \
754  vec_t Ret = sv; \
755  sv OPASSIGN static_cast<DataT>(1); \
756  return Ret; \
757  }
758 
759  __SYCL_UOP(++, +=)
760  __SYCL_UOP(--, -=)
761 #undef __SYCL_UOP
762 
763  template <typename T = DataT>
764  friend typename std::enable_if_t<
765  std::is_same_v<T, DataT> && !detail::is_vgenfloat_v<T>, vec_t>
766  operator~(const SwizzleOp &Rhs) {
767  vec_t Tmp = Rhs;
768  return ~Tmp;
769  }
770 
771  friend vec_rel_t operator!(const SwizzleOp &Rhs) {
772  vec_t Tmp = Rhs;
773  return !Tmp;
774  }
775 
776  friend vec_t operator+(const SwizzleOp &Rhs) {
777  vec_t Tmp = Rhs;
778  return +Tmp;
779  }
780 
781  friend vec_t operator-(const SwizzleOp &Rhs) {
782  vec_t Tmp = Rhs;
783  return -Tmp;
784  }
785 
786 // scalar BINOP vec<>
787 // scalar BINOP SwizzleOp
788 // vec<> BINOP SwizzleOp
789 #ifdef __SYCL_BINOP
790 #error "Undefine __SYCL_BINOP macro"
791 #endif
792 #define __SYCL_BINOP(BINOP, COND) \
793  template <typename T = DataT> \
794  friend std::enable_if_t<(COND), vec_t> operator BINOP( \
795  const DataT & Lhs, const SwizzleOp & Rhs) { \
796  vec_t Tmp = Rhs; \
797  return Lhs BINOP Tmp; \
798  } \
799  template <typename T = DataT> \
800  friend std::enable_if_t<(COND), vec_t> operator BINOP(const SwizzleOp & Lhs, \
801  const DataT & Rhs) { \
802  vec_t Tmp = Lhs; \
803  return Tmp BINOP Rhs; \
804  } \
805  template <typename T = DataT> \
806  friend std::enable_if_t<(COND), vec_t> operator BINOP( \
807  const vec_t & Lhs, const SwizzleOp & Rhs) { \
808  vec_t Tmp = Rhs; \
809  return Lhs BINOP Tmp; \
810  } \
811  template <typename T = DataT> \
812  friend std::enable_if_t<(COND), vec_t> operator BINOP(const SwizzleOp & Lhs, \
813  const vec_t & Rhs) { \
814  vec_t Tmp = Lhs; \
815  return Tmp BINOP Rhs; \
816  }
817 
818  __SYCL_BINOP(+, (!detail::is_byte_v<T>))
819  __SYCL_BINOP(-, (!detail::is_byte_v<T>))
820  __SYCL_BINOP(*, (!detail::is_byte_v<T>))
821  __SYCL_BINOP(/, (!detail::is_byte_v<T>))
822  __SYCL_BINOP(%, (!detail::is_byte_v<T>))
823  __SYCL_BINOP(&, true)
824  __SYCL_BINOP(|, true)
825  __SYCL_BINOP(^, true)
826  // We have special <<, >> operators for std::byte.
827  __SYCL_BINOP(>>, (!detail::is_byte_v<T>))
828  __SYCL_BINOP(<<, (!detail::is_byte_v<T>))
829 
830  template <typename T = DataT>
831  friend std::enable_if_t<detail::is_byte_v<T>, vec_t>
832  operator>>(const SwizzleOp &Lhs, const int shift) {
833  vec_t Tmp = Lhs;
834  return Tmp >> shift;
835  }
836 
837  template <typename T = DataT>
838  friend std::enable_if_t<detail::is_byte_v<T>, vec_t>
839  operator<<(const SwizzleOp &Lhs, const int shift) {
840  vec_t Tmp = Lhs;
841  return Tmp << shift;
842  }
843 #undef __SYCL_BINOP
844 
845 // scalar RELLOGOP vec<>
846 // scalar RELLOGOP SwizzleOp
847 // vec<> RELLOGOP SwizzleOp
848 #ifdef __SYCL_RELLOGOP
849 #error "Undefine __SYCL_RELLOGOP macro"
850 #endif
851 #define __SYCL_RELLOGOP(RELLOGOP, COND) \
852  template <typename T = DataT> \
853  friend std::enable_if_t<(COND), vec_rel_t> operator RELLOGOP( \
854  const DataT & Lhs, const SwizzleOp & Rhs) { \
855  vec_t Tmp = Rhs; \
856  return Lhs RELLOGOP Tmp; \
857  } \
858  template <typename T = DataT> \
859  friend std::enable_if_t<(COND), vec_rel_t> operator RELLOGOP( \
860  const SwizzleOp & Lhs, const DataT & Rhs) { \
861  vec_t Tmp = Lhs; \
862  return Tmp RELLOGOP Rhs; \
863  } \
864  template <typename T = DataT> \
865  friend std::enable_if_t<(COND), vec_rel_t> operator RELLOGOP( \
866  const vec_t & Lhs, const SwizzleOp & Rhs) { \
867  vec_t Tmp = Rhs; \
868  return Lhs RELLOGOP Tmp; \
869  } \
870  template <typename T = DataT> \
871  friend std::enable_if_t<(COND), vec_rel_t> operator RELLOGOP( \
872  const SwizzleOp & Lhs, const vec_t & Rhs) { \
873  vec_t Tmp = Lhs; \
874  return Tmp RELLOGOP Rhs; \
875  }
876 
877  __SYCL_RELLOGOP(==, (!detail::is_byte_v<T>))
878  __SYCL_RELLOGOP(!=, (!detail::is_byte_v<T>))
879  __SYCL_RELLOGOP(>, (!detail::is_byte_v<T>))
880  __SYCL_RELLOGOP(<, (!detail::is_byte_v<T>))
881  __SYCL_RELLOGOP(>=, (!detail::is_byte_v<T>))
882  __SYCL_RELLOGOP(<=, (!detail::is_byte_v<T>))
883  __SYCL_RELLOGOP(&&, (!detail::is_byte_v<T> && !detail::is_vgenfloat_v<T>))
884  __SYCL_RELLOGOP(||, (!detail::is_byte_v<T> && !detail::is_vgenfloat_v<T>))
885 #undef __SYCL_RELLOGOP
886 
887  template <int IdxNum = getNumElements(),
888  typename = EnableIfMultipleIndexes<IdxNum>>
889  SwizzleOp &operator=(const vec<DataT, IdxNum> &Rhs) {
890  std::array<int, IdxNum> Idxs{Indexes...};
891  for (size_t I = 0; I < Idxs.size(); ++I) {
892  (*m_Vector)[Idxs[I]] = Rhs[I];
893  }
894  return *this;
895  }
896 
897  template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
898  SwizzleOp &operator=(const DataT &Rhs) {
899  std::array<int, IdxNum> Idxs{Indexes...};
900  (*m_Vector)[Idxs[0]] = Rhs;
901  return *this;
902  }
903 
904  template <int IdxNum = getNumElements(),
905  EnableIfMultipleIndexes<IdxNum, bool> = true>
906  SwizzleOp &operator=(const DataT &Rhs) {
907  std::array<int, IdxNum> Idxs{Indexes...};
908  for (auto Idx : Idxs) {
909  (*m_Vector)[Idx] = Rhs;
910  }
911  return *this;
912  }
913 
914  template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
915  SwizzleOp &operator=(DataT &&Rhs) {
916  std::array<int, IdxNum> Idxs{Indexes...};
917  (*m_Vector)[Idxs[0]] = Rhs;
918  return *this;
919  }
920 
921  template <typename T, typename = EnableIfScalarType<T>>
922  NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>
923  operator*(const T &Rhs) const {
924  return NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
925  m_Vector, *this, GetScalarOp<T>(Rhs));
926  }
927 
928  template <typename RhsOperation,
929  typename = EnableIfNoScalarType<RhsOperation>>
930  NewLHOp<RhsOperation, std::multiplies, Indexes...>
931  operator*(const RhsOperation &Rhs) const {
932  return NewLHOp<RhsOperation, std::multiplies, Indexes...>(m_Vector, *this,
933  Rhs);
934  }
935 
936  template <typename T, typename = EnableIfScalarType<T>>
937  NewLHOp<GetScalarOp<T>, std::plus, Indexes...> operator+(const T &Rhs) const {
938  return NewLHOp<GetScalarOp<T>, std::plus, Indexes...>(m_Vector, *this,
939  GetScalarOp<T>(Rhs));
940  }
941 
942  template <typename RhsOperation,
943  typename = EnableIfNoScalarType<RhsOperation>>
944  NewLHOp<RhsOperation, std::plus, Indexes...>
945  operator+(const RhsOperation &Rhs) const {
946  return NewLHOp<RhsOperation, std::plus, Indexes...>(m_Vector, *this, Rhs);
947  }
948 
949  template <typename T, typename = EnableIfScalarType<T>>
950  NewLHOp<GetScalarOp<T>, std::minus, Indexes...>
951  operator-(const T &Rhs) const {
952  return NewLHOp<GetScalarOp<T>, std::minus, Indexes...>(m_Vector, *this,
953  GetScalarOp<T>(Rhs));
954  }
955 
956  template <typename RhsOperation,
957  typename = EnableIfNoScalarType<RhsOperation>>
958  NewLHOp<RhsOperation, std::minus, Indexes...>
959  operator-(const RhsOperation &Rhs) const {
960  return NewLHOp<RhsOperation, std::minus, Indexes...>(m_Vector, *this, Rhs);
961  }
962 
963  template <typename T, typename = EnableIfScalarType<T>>
964  NewLHOp<GetScalarOp<T>, std::divides, Indexes...>
965  operator/(const T &Rhs) const {
966  return NewLHOp<GetScalarOp<T>, std::divides, Indexes...>(
967  m_Vector, *this, GetScalarOp<T>(Rhs));
968  }
969 
970  template <typename RhsOperation,
971  typename = EnableIfNoScalarType<RhsOperation>>
972  NewLHOp<RhsOperation, std::divides, Indexes...>
973  operator/(const RhsOperation &Rhs) const {
974  return NewLHOp<RhsOperation, std::divides, Indexes...>(m_Vector, *this,
975  Rhs);
976  }
977 
978  template <typename T, typename = EnableIfScalarType<T>>
979  NewLHOp<GetScalarOp<T>, std::modulus, Indexes...>
980  operator%(const T &Rhs) const {
981  return NewLHOp<GetScalarOp<T>, std::modulus, Indexes...>(
982  m_Vector, *this, GetScalarOp<T>(Rhs));
983  }
984 
985  template <typename RhsOperation,
986  typename = EnableIfNoScalarType<RhsOperation>>
987  NewLHOp<RhsOperation, std::modulus, Indexes...>
988  operator%(const RhsOperation &Rhs) const {
989  return NewLHOp<RhsOperation, std::modulus, Indexes...>(m_Vector, *this,
990  Rhs);
991  }
992 
993  template <typename T, typename = EnableIfScalarType<T>>
994  NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>
995  operator&(const T &Rhs) const {
996  return NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>(
997  m_Vector, *this, GetScalarOp<T>(Rhs));
998  }
999 
1000  template <typename RhsOperation,
1001  typename = EnableIfNoScalarType<RhsOperation>>
1002  NewLHOp<RhsOperation, std::bit_and, Indexes...>
1003  operator&(const RhsOperation &Rhs) const {
1004  return NewLHOp<RhsOperation, std::bit_and, Indexes...>(m_Vector, *this,
1005  Rhs);
1006  }
1007 
1008  template <typename T, typename = EnableIfScalarType<T>>
1009  NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>
1010  operator|(const T &Rhs) const {
1011  return NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>(
1012  m_Vector, *this, GetScalarOp<T>(Rhs));
1013  }
1014 
1015  template <typename RhsOperation,
1016  typename = EnableIfNoScalarType<RhsOperation>>
1017  NewLHOp<RhsOperation, std::bit_or, Indexes...>
1018  operator|(const RhsOperation &Rhs) const {
1019  return NewLHOp<RhsOperation, std::bit_or, Indexes...>(m_Vector, *this, Rhs);
1020  }
1021 
1022  template <typename T, typename = EnableIfScalarType<T>>
1023  NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>
1024  operator^(const T &Rhs) const {
1025  return NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>(
1026  m_Vector, *this, GetScalarOp<T>(Rhs));
1027  }
1028 
1029  template <typename RhsOperation,
1030  typename = EnableIfNoScalarType<RhsOperation>>
1031  NewLHOp<RhsOperation, std::bit_xor, Indexes...>
1032  operator^(const RhsOperation &Rhs) const {
1033  return NewLHOp<RhsOperation, std::bit_xor, Indexes...>(m_Vector, *this,
1034  Rhs);
1035  }
1036 
1037  template <typename T, typename = EnableIfScalarType<T>>
1038  NewLHOp<GetScalarOp<T>, RShift, Indexes...> operator>>(const T &Rhs) const {
1039  return NewLHOp<GetScalarOp<T>, RShift, Indexes...>(m_Vector, *this,
1040  GetScalarOp<T>(Rhs));
1041  }
1042 
1043  template <typename RhsOperation,
1044  typename = EnableIfNoScalarType<RhsOperation>>
1045  NewLHOp<RhsOperation, RShift, Indexes...>
1046  operator>>(const RhsOperation &Rhs) const {
1047  return NewLHOp<RhsOperation, RShift, Indexes...>(m_Vector, *this, Rhs);
1048  }
1049 
1050  template <typename T, typename = EnableIfScalarType<T>>
1051  NewLHOp<GetScalarOp<T>, LShift, Indexes...> operator<<(const T &Rhs) const {
1052  return NewLHOp<GetScalarOp<T>, LShift, Indexes...>(m_Vector, *this,
1053  GetScalarOp<T>(Rhs));
1054  }
1055 
1056  template <typename RhsOperation,
1057  typename = EnableIfNoScalarType<RhsOperation>>
1058  NewLHOp<RhsOperation, LShift, Indexes...>
1059  operator<<(const RhsOperation &Rhs) const {
1060  return NewLHOp<RhsOperation, LShift, Indexes...>(m_Vector, *this, Rhs);
1061  }
1062 
1063  template <
1064  typename T1, typename T2, typename T3, template <typename> class T4,
1065  int... T5,
1066  typename = typename std::enable_if_t<sizeof...(T5) == getNumElements()>>
1067  SwizzleOp &operator=(const SwizzleOp<T1, T2, T3, T4, T5...> &Rhs) {
1068  std::array<int, getNumElements()> Idxs{Indexes...};
1069  for (size_t I = 0; I < Idxs.size(); ++I) {
1070  (*m_Vector)[Idxs[I]] = Rhs.getValue(I);
1071  }
1072  return *this;
1073  }
1074 
1075  template <
1076  typename T1, typename T2, typename T3, template <typename> class T4,
1077  int... T5,
1078  typename = typename std::enable_if_t<sizeof...(T5) == getNumElements()>>
1079  SwizzleOp &operator=(SwizzleOp<T1, T2, T3, T4, T5...> &&Rhs) {
1080  std::array<int, getNumElements()> Idxs{Indexes...};
1081  for (size_t I = 0; I < Idxs.size(); ++I) {
1082  (*m_Vector)[Idxs[I]] = Rhs.getValue(I);
1083  }
1084  return *this;
1085  }
1086 
1087  template <typename T, typename = EnableIfScalarType<T>>
1088  NewRelOp<GetScalarOp<T>, EqualTo, Indexes...> operator==(const T &Rhs) const {
1089  return NewRelOp<GetScalarOp<T>, EqualTo, Indexes...>(NULL, *this,
1090  GetScalarOp<T>(Rhs));
1091  }
1092 
1093  template <typename RhsOperation,
1094  typename = EnableIfNoScalarType<RhsOperation>>
1095  NewRelOp<RhsOperation, EqualTo, Indexes...>
1096  operator==(const RhsOperation &Rhs) const {
1097  return NewRelOp<RhsOperation, EqualTo, Indexes...>(NULL, *this, Rhs);
1098  }
1099 
1100  template <typename T, typename = EnableIfScalarType<T>>
1101  NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>
1102  operator!=(const T &Rhs) const {
1103  return NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>(
1104  NULL, *this, GetScalarOp<T>(Rhs));
1105  }
1106 
1107  template <typename RhsOperation,
1108  typename = EnableIfNoScalarType<RhsOperation>>
1109  NewRelOp<RhsOperation, NotEqualTo, Indexes...>
1110  operator!=(const RhsOperation &Rhs) const {
1111  return NewRelOp<RhsOperation, NotEqualTo, Indexes...>(NULL, *this, Rhs);
1112  }
1113 
1114  template <typename T, typename = EnableIfScalarType<T>>
1115  NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>
1116  operator>=(const T &Rhs) const {
1117  return NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>(
1118  NULL, *this, GetScalarOp<T>(Rhs));
1119  }
1120 
1121  template <typename RhsOperation,
1122  typename = EnableIfNoScalarType<RhsOperation>>
1123  NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>
1124  operator>=(const RhsOperation &Rhs) const {
1125  return NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>(NULL, *this, Rhs);
1126  }
1127 
1128  template <typename T, typename = EnableIfScalarType<T>>
1129  NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>
1130  operator<=(const T &Rhs) const {
1131  return NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>(
1132  NULL, *this, GetScalarOp<T>(Rhs));
1133  }
1134 
1135  template <typename RhsOperation,
1136  typename = EnableIfNoScalarType<RhsOperation>>
1137  NewRelOp<RhsOperation, LessEqualTo, Indexes...>
1138  operator<=(const RhsOperation &Rhs) const {
1139  return NewRelOp<RhsOperation, LessEqualTo, Indexes...>(NULL, *this, Rhs);
1140  }
1141 
1142  template <typename T, typename = EnableIfScalarType<T>>
1143  NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>
1144  operator>(const T &Rhs) const {
1145  return NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>(
1146  NULL, *this, GetScalarOp<T>(Rhs));
1147  }
1148 
1149  template <typename RhsOperation,
1150  typename = EnableIfNoScalarType<RhsOperation>>
1151  NewRelOp<RhsOperation, GreaterThan, Indexes...>
1152  operator>(const RhsOperation &Rhs) const {
1153  return NewRelOp<RhsOperation, GreaterThan, Indexes...>(NULL, *this, Rhs);
1154  }
1155 
1156  template <typename T, typename = EnableIfScalarType<T>>
1157  NewRelOp<GetScalarOp<T>, LessThan, Indexes...> operator<(const T &Rhs) const {
1158  return NewRelOp<GetScalarOp<T>, LessThan, Indexes...>(NULL, *this,
1159  GetScalarOp<T>(Rhs));
1160  }
1161 
1162  template <typename RhsOperation,
1163  typename = EnableIfNoScalarType<RhsOperation>>
1164  NewRelOp<RhsOperation, LessThan, Indexes...>
1165  operator<(const RhsOperation &Rhs) const {
1166  return NewRelOp<RhsOperation, LessThan, Indexes...>(NULL, *this, Rhs);
1167  }
1168 
1169  template <typename T, typename = EnableIfScalarType<T>>
1170  NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>
1171  operator&&(const T &Rhs) const {
1172  return NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>(
1173  NULL, *this, GetScalarOp<T>(Rhs));
1174  }
1175 
1176  template <typename RhsOperation,
1177  typename = EnableIfNoScalarType<RhsOperation>>
1178  NewRelOp<RhsOperation, LogicalAnd, Indexes...>
1179  operator&&(const RhsOperation &Rhs) const {
1180  return NewRelOp<RhsOperation, LogicalAnd, Indexes...>(NULL, *this, Rhs);
1181  }
1182 
1183  template <typename T, typename = EnableIfScalarType<T>>
1184  NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>
1185  operator||(const T &Rhs) const {
1186  return NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>(NULL, *this,
1187  GetScalarOp<T>(Rhs));
1188  }
1189 
1190  template <typename RhsOperation,
1191  typename = EnableIfNoScalarType<RhsOperation>>
1192  NewRelOp<RhsOperation, LogicalOr, Indexes...>
1193  operator||(const RhsOperation &Rhs) const {
1194  return NewRelOp<RhsOperation, LogicalOr, Indexes...>(NULL, *this, Rhs);
1195  }
1196 
1197  // Begin hi/lo, even/odd, xyzw, and rgba swizzles.
1198 private:
1199  // Indexer used in the swizzles.def.
1200  // Currently it is defined as a template struct. Replacing it with a constexpr
1201  // function would activate a bug in MSVC that is fixed only in v19.20.
1202  // Until then MSVC does not recognize such constexpr functions as const and
1203  // thus does not let using them in template parameters inside swizzle.def.
1204  template <int Index> struct Indexer {
1205  static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...};
1206  static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index];
1207  };
1208 
1209 public:
1210 #ifdef __SYCL_ACCESS_RETURN
1211 #error "Undefine __SYCL_ACCESS_RETURN macro"
1212 #endif
1213 #define __SYCL_ACCESS_RETURN m_Vector
1214 #include "swizzles.def"
1215 #undef __SYCL_ACCESS_RETURN
1216  // End of hi/lo, even/odd, xyzw, and rgba swizzles.
1217 
1218  // Leave store() interface to automatic conversion to vec<>.
1219  // Load to vec_t and then assign to swizzle.
1220  template <access::address_space Space, access::decorated DecorateAddress>
1221  void load(size_t offset, multi_ptr<DataT, Space, DecorateAddress> ptr) {
1222  vec_t Tmp;
1223  Tmp.load(offset, ptr);
1224  *this = Tmp;
1225  }
1226 
1227  template <typename convertT, rounding_mode roundingMode>
1228  vec<convertT, sizeof...(Indexes)> convert() const {
1229  // First materialize the swizzle to vec_t and then apply convert() to it.
1230  vec_t Tmp;
1231  std::array<int, getNumElements()> Idxs{Indexes...};
1232  for (size_t I = 0; I < Idxs.size(); ++I) {
1233  Tmp[I] = (*m_Vector)[Idxs[I]];
1234  }
1235  return Tmp.template convert<convertT, roundingMode>();
1236  }
1237 
1238  template <typename asT> asT as() const {
1239  // First materialize the swizzle to vec_t and then apply as() to it.
1240  vec_t Tmp = *this;
1241  static_assert((sizeof(Tmp) == sizeof(asT)),
1242  "The new SYCL vec type must have the same storage size in "
1243  "bytes as this SYCL swizzled vec");
1244  static_assert(
1245  detail::is_contained<asT, detail::gtl::vector_basic_list>::value ||
1246  detail::is_contained<asT, detail::gtl::vector_bool_list>::value,
1247  "asT must be SYCL vec of a different element type and "
1248  "number of elements specified by asT");
1249  return Tmp.template as<asT>();
1250  }
1251 
1252 private:
1253  SwizzleOp(const SwizzleOp &Rhs)
1254  : m_Vector(Rhs.m_Vector), m_LeftOperation(Rhs.m_LeftOperation),
1255  m_RightOperation(Rhs.m_RightOperation) {}
1256 
1257  SwizzleOp(VecT *Vector, OperationLeftT LeftOperation,
1258  OperationRightT RightOperation)
1259  : m_Vector(Vector), m_LeftOperation(LeftOperation),
1260  m_RightOperation(RightOperation) {}
1261 
1262  SwizzleOp(VecT *Vector) : m_Vector(Vector) {}
1263 
1264  SwizzleOp(SwizzleOp &&Rhs)
1265  : m_Vector(Rhs.m_Vector), m_LeftOperation(std::move(Rhs.m_LeftOperation)),
1266  m_RightOperation(std::move(Rhs.m_RightOperation)) {}
1267 
1268  // Either performing CurrentOperation on results of left and right operands
1269  // or reading values from actual vector. Perform implicit type conversion when
1270  // the number of elements == 1
1271 
1272  template <int IdxNum = getNumElements()>
1273  CommonDataT getValue(EnableIfOneIndex<IdxNum, size_t> Index) const {
1274  if (std::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
1275  std::array<int, getNumElements()> Idxs{Indexes...};
1276  return (*m_Vector)[Idxs[Index]];
1277  }
1278  auto Op = OperationCurrentT<CommonDataT>();
1279  return Op(m_LeftOperation.getValue(Index),
1280  m_RightOperation.getValue(Index));
1281  }
1282 
1283  template <int IdxNum = getNumElements()>
1284  DataT getValue(EnableIfMultipleIndexes<IdxNum, size_t> Index) const {
1285  if (std::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
1286  std::array<int, getNumElements()> Idxs{Indexes...};
1287  return (*m_Vector)[Idxs[Index]];
1288  }
1289  auto Op = OperationCurrentT<DataT>();
1290  return Op(m_LeftOperation.getValue(Index),
1291  m_RightOperation.getValue(Index));
1292  }
1293 
1294  template <template <typename> class Operation, typename RhsOperation>
1295  void operatorHelper(const RhsOperation &Rhs) const {
1296  Operation<DataT> Op;
1297  std::array<int, getNumElements()> Idxs{Indexes...};
1298  for (size_t I = 0; I < Idxs.size(); ++I) {
1299  DataT Res = Op((*m_Vector)[Idxs[I]], Rhs.getValue(I));
1300  (*m_Vector)[Idxs[I]] = Res;
1301  }
1302  }
1303 
1304  // fields
1305  VecT *m_Vector;
1306 
1307  OperationLeftT m_LeftOperation;
1308  OperationRightT m_RightOperation;
1309 
1310  // friends
1311  template <typename T1, int T2> friend class sycl::vec;
1312 
1313  template <typename T1, typename T2, typename T3, template <typename> class T4,
1314  int... T5>
1315  friend class SwizzleOp;
1316 };
1318 } // namespace detail
1319 } // namespace _V1
1320 } // namespace sycl
sycl::_V1::detail::GetOp
Definition: vector.hpp:95
sycl::_V1::detail::GetOp::DataT
T DataT
Definition: vector.hpp:97
sycl::_V1::detail::GetOp::operator()
DataT operator()(DataT, DataT)
Definition: vector.hpp:99
sycl::_V1::detail::GetOp::getValue
DataT getValue(size_t) const
Definition: vector.hpp:98
sycl::_V1::detail::vec_arith
Definition: vector_arith.hpp:113
sycl::_V1::vec
Definition: vector.hpp:133
defines_elementary.hpp
__SYCL_EBO
#define __SYCL_EBO
Definition: defines_elementary.hpp:43
generic_type_lists.hpp
generic_type_traits.hpp
half_type.hpp
sycl::_V1::detail::half_impl::StorageT
uint16_t StorageT
Definition: half_type.hpp:172
sycl::_V1::detail::bfloat16
sycl::ext::oneapi::bfloat16 bfloat16
Definition: vector_convert.hpp:217
sycl::_V1::detail::ConvertToOpenCLType_t
decltype(convertToOpenCLType(std::declval< T >())) ConvertToOpenCLType_t
Definition: generic_type_traits.hpp:387
sycl::_V1::ext::intel::experimental::byte
unsigned char byte
Definition: online_compiler.hpp:22
sycl::_V1::ext::oneapi::detail::Bfloat16StorageT
uint16_t Bfloat16StorageT
Definition: bfloat16.hpp:53
sycl::_V1::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:464
sycl::_V1::half
sycl::detail::half_impl::half half
Definition: aliases.hpp:101
sycl
Definition: access.hpp:18
sycl::rounding_mode
rounding_mode
Definition: vector.hpp:57
sycl::_V1::detail::ScalarConversionOperatorMixIn
Definition: vector.hpp:117
sycl::_V1::detail::map_type
Definition: type_traits.hpp:553
sycl::_V1::elem
Definition: vector.hpp:61
sycl::_V1::elem::g
static constexpr int g
Definition: vector.hpp:67
sycl::_V1::elem::s4
static constexpr int s4
Definition: vector.hpp:74
sycl::_V1::elem::s9
static constexpr int s9
Definition: vector.hpp:79
sycl::_V1::elem::s5
static constexpr int s5
Definition: vector.hpp:75
sycl::_V1::elem::s1
static constexpr int s1
Definition: vector.hpp:71
sycl::_V1::elem::s3
static constexpr int s3
Definition: vector.hpp:73
sycl::_V1::elem::s0
static constexpr int s0
Definition: vector.hpp:70
sycl::_V1::elem::sF
static constexpr int sF
Definition: vector.hpp:85
sycl::_V1::elem::s7
static constexpr int s7
Definition: vector.hpp:77
sycl::_V1::elem::x
static constexpr int x
Definition: vector.hpp:62
sycl::_V1::elem::z
static constexpr int z
Definition: vector.hpp:64
sycl::_V1::elem::s8
static constexpr int s8
Definition: vector.hpp:78
sycl::_V1::elem::a
static constexpr int a
Definition: vector.hpp:69
sycl::_V1::elem::sD
static constexpr int sD
Definition: vector.hpp:83
sycl::_V1::elem::y
static constexpr int y
Definition: vector.hpp:63
sycl::_V1::elem::w
static constexpr int w
Definition: vector.hpp:65
sycl::_V1::elem::sC
static constexpr int sC
Definition: vector.hpp:82
sycl::_V1::elem::sB
static constexpr int sB
Definition: vector.hpp:81
sycl::_V1::elem::sA
static constexpr int sA
Definition: vector.hpp:80
sycl::_V1::elem::s2
static constexpr int s2
Definition: vector.hpp:72
sycl::_V1::elem::b
static constexpr int b
Definition: vector.hpp:68
sycl::_V1::elem::r
static constexpr int r
Definition: vector.hpp:66
sycl::_V1::elem::sE
static constexpr int sE
Definition: vector.hpp:84
sycl::_V1::elem::s6
static constexpr int s6
Definition: vector.hpp:76
type_list.hpp
type_traits.hpp
vector_arith.hpp