DPC++ Runtime
Runtime libraries for oneAPI DPC++
types.hpp
Go to the documentation of this file.
1 //==---------------- types.hpp --- SYCL types ------------------------------==//
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 // Implements vec and __swizzled_vec__ classes, this older implementation
10 // is not fully conforming, but is used by default. Use
11 // -fpreview-breaking-changes to use the newer conforming version.
12 
13 #pragma once
14 
15 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
16 #include <sycl/detail/generic_type_traits.hpp>
17 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
18 
19 // Check if Clang's ext_vector_type attribute is available. Host compiler
20 // may not be Clang, and Clang may not be built with the extension.
21 #ifdef __clang__
22 #ifndef __has_extension
23 #define __has_extension(x) 0
24 #endif
25 #ifdef __HAS_EXT_VECTOR_TYPE__
26 #error "Undefine __HAS_EXT_VECTOR_TYPE__ macro"
27 #endif
28 #if __has_extension(attribute_ext_vector_type)
29 #define __HAS_EXT_VECTOR_TYPE__
30 #endif
31 #endif // __clang__
32 
33 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
34 
35 #if !defined(__HAS_EXT_VECTOR_TYPE__) && defined(__SYCL_DEVICE_ONLY__)
36 #error "SYCL device compiler is built without ext_vector_type support"
37 #endif
38 
39 #if defined(__SYCL_DEVICE_ONLY__)
40 #define __SYCL_USE_EXT_VECTOR_TYPE__
41 #endif
42 
43 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
44 
45 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
46 
47 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
48 #error "Undefine __SYCL_USE_EXT_VECTOR_TYPE__ macro"
49 #endif
50 #ifdef __HAS_EXT_VECTOR_TYPE__
51 #if defined(__SYCL_DEVICE_ONLY__) || !defined(__NO_EXT_VECTOR_TYPE_ON_HOST__)
52 #define __SYCL_USE_EXT_VECTOR_TYPE__
53 #endif
54 #elif defined(__SYCL_DEVICE_ONLY__)
55 // This is a soft error. We expect the device compiler to have ext_vector_type
56 // support, but that should not be a hard requirement.
57 #error "SYCL device compiler is built without ext_vector_type support"
58 #endif // __HAS_EXT_VECTOR_TYPE__
59 
60 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
61 
62 #include <sycl/access/access.hpp> // for decorated, address_space
63 #include <sycl/aliases.hpp> // for half, cl_char, cl_int
64 #include <sycl/detail/common.hpp> // for ArrayCreator, RepeatV...
65 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
66 #include <sycl/detail/generic_type_lists.hpp> // for vector_basic_list
67 #include <sycl/detail/generic_type_traits.hpp> // for is_sigeninteger, is_s...
68 #include <sycl/detail/memcpy.hpp> // for memcpy
69 #include <sycl/detail/type_list.hpp> // for is_contained
70 #include <sycl/detail/type_traits.hpp> // for is_floating_point
71 #include <sycl/detail/vector_convert.hpp> // for convertImpl
72 #include <sycl/detail/vector_traits.hpp> // for vector_alignment
73 #include <sycl/exception.hpp> // for make_error_code, errc
74 #include <sycl/half_type.hpp> // for StorageT, half, Vec16...
75 #include <sycl/marray.hpp> // for __SYCL_BINOP, __SYCL_...
76 #include <sycl/multi_ptr.hpp> // for multi_ptr
77 
78 #include <array> // for array
79 #include <assert.h> // for assert
80 #include <cstddef> // for size_t, NULL, byte
81 #include <cstdint> // for uint8_t, int16_t, int...
82 #include <functional> // for divides, multiplies
83 #include <iterator> // for pair
84 #include <optional> // for optional
85 #include <ostream> // for operator<<, basic_ost...
86 #include <tuple> // for tuple
87 #include <type_traits> // for enable_if_t, is_same
88 #include <utility> // for index_sequence, make_...
89 #include <variant> // for tuple, variant
90 
91 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
92 
93 #ifndef __SYCL_USE_EXT_VECTOR_TYPE__
94 #include <sycl/detail/cl.h>
95 #endif
96 #include <cstring>
97 
98 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
99 
100 // 4.10.1: Scalar data types
101 // 4.10.2: SYCL vector types
102 
103 namespace sycl {
104 inline namespace _V1 {
105 
106 struct elem {
107  static constexpr int x = 0;
108  static constexpr int y = 1;
109  static constexpr int z = 2;
110  static constexpr int w = 3;
111  static constexpr int r = 0;
112  static constexpr int g = 1;
113  static constexpr int b = 2;
114  static constexpr int a = 3;
115  static constexpr int s0 = 0;
116  static constexpr int s1 = 1;
117  static constexpr int s2 = 2;
118  static constexpr int s3 = 3;
119  static constexpr int s4 = 4;
120  static constexpr int s5 = 5;
121  static constexpr int s6 = 6;
122  static constexpr int s7 = 7;
123  static constexpr int s8 = 8;
124  static constexpr int s9 = 9;
125  static constexpr int sA = 10;
126  static constexpr int sB = 11;
127  static constexpr int sC = 12;
128  static constexpr int sD = 13;
129  static constexpr int sE = 14;
130  static constexpr int sF = 15;
131 };
132 
133 namespace detail {
134 // select_apply_cl_t selects from T8/T16/T32/T64 basing on
135 // sizeof(_IN). expected to handle scalar types in _IN.
136 template <typename _IN, typename T8, typename T16, typename T32, typename T64>
137 using select_apply_cl_t = std::conditional_t<
138  sizeof(_IN) == 1, T8,
139  std::conditional_t<sizeof(_IN) == 2, T16,
140  std::conditional_t<sizeof(_IN) == 4, T32, T64>>>;
141 
142 template <typename T> struct vec_helper {
143  using RetType = T;
144  static constexpr RetType get(T value) { return value; }
145 };
146 template <> struct vec_helper<bool> {
147  using RetType = select_apply_cl_t<bool, std::int8_t, std::int16_t,
148  std::int32_t, std::int64_t>;
149  static constexpr RetType get(bool value) { return value; }
150 };
151 
152 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
153 template <> struct vec_helper<std::byte> {
154  using RetType = std::uint8_t;
155  static constexpr RetType get(std::byte value) { return (RetType)value; }
156  static constexpr std::byte get(std::uint8_t value) {
157  return (std::byte)value;
158  }
159 };
160 #endif
161 
162 template <typename VecT, typename OperationLeftT, typename OperationRightT,
163  template <typename> class OperationCurrentT, int... Indexes>
164 class SwizzleOp;
165 
166 template <typename T, int N, typename V = void> struct VecStorage;
167 
168 // Element type for relational operator return value.
169 template <typename DataT>
170 using rel_t = typename std::conditional_t<
171  sizeof(DataT) == sizeof(opencl::cl_char), opencl::cl_char,
172  typename std::conditional_t<
173  sizeof(DataT) == sizeof(opencl::cl_short), opencl::cl_short,
174  typename std::conditional_t<
175  sizeof(DataT) == sizeof(opencl::cl_int), opencl::cl_int,
176  typename std::conditional_t<sizeof(DataT) ==
177  sizeof(opencl::cl_long),
178  opencl::cl_long, bool>>>>;
179 
180 // Special type indicating that SwizzleOp should just read value from vector -
181 // not trying to perform any operations. Should not be called.
182 template <typename T> class GetOp {
183 public:
184  using DataT = T;
185  DataT getValue(size_t) const { return (DataT)0; }
186  DataT operator()(DataT, DataT) { return (DataT)0; }
187 };
188 
189 // Special type for working SwizzleOp with scalars, stores a scalar and gives
190 // the scalar at any index. Provides interface is compatible with SwizzleOp
191 // operations
192 template <typename T> class GetScalarOp {
193 public:
194  using DataT = T;
195  GetScalarOp(DataT Data) : m_Data(Data) {}
196  DataT getValue(size_t) const { return m_Data; }
197 
198 private:
199  DataT m_Data;
200 };
201 
202 template <typename T> struct EqualTo {
203  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
204  return (Lhs == Rhs) ? -1 : 0;
205  }
206 };
207 
208 template <typename T> struct NotEqualTo {
209  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
210  return (Lhs != Rhs) ? -1 : 0;
211  }
212 };
213 
214 template <typename T> struct GreaterEqualTo {
215  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
216  return (Lhs >= Rhs) ? -1 : 0;
217  }
218 };
219 
220 template <typename T> struct LessEqualTo {
221  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
222  return (Lhs <= Rhs) ? -1 : 0;
223  }
224 };
225 
226 template <typename T> struct GreaterThan {
227  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
228  return (Lhs > Rhs) ? -1 : 0;
229  }
230 };
231 
232 template <typename T> struct LessThan {
233  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
234  return (Lhs < Rhs) ? -1 : 0;
235  }
236 };
237 
238 template <typename T> struct LogicalAnd {
239  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
240  return (Lhs && Rhs) ? -1 : 0;
241  }
242 };
243 
244 template <typename T> struct LogicalOr {
245  constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {
246  return (Lhs || Rhs) ? -1 : 0;
247  }
248 };
249 
250 template <typename T> struct RShift {
251  constexpr T operator()(const T &Lhs, const T &Rhs) const {
252  return Lhs >> Rhs;
253  }
254 };
255 
256 template <typename T> struct LShift {
257  constexpr T operator()(const T &Lhs, const T &Rhs) const {
258  return Lhs << Rhs;
259  }
260 };
261 
262 // Forward declarations
263 template <typename TransformedArgType, int Dims, typename KernelType>
264 class RoundedRangeKernel;
265 template <typename TransformedArgType, int Dims, typename KernelType>
266 class RoundedRangeKernelWithKH;
267 
268 } // namespace detail
269 
270 template <typename T> using vec_data = detail::vec_helper<T>;
271 
272 template <typename T>
273 using vec_data_t = typename detail::vec_helper<T>::RetType;
274 
275 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
276 
277 #if defined(_WIN32) && (_MSC_VER)
278 // MSVC Compiler doesn't allow using of function arguments with alignment
279 // requirements. MSVC Compiler Error C2719: 'parameter': formal parameter with
280 // __declspec(align('#')) won't be aligned. The align __declspec modifier
281 // is not permitted on function parameters. Function parameter alignment
282 // is controlled by the calling convention used.
283 // For more information, see Calling Conventions
284 // (https://docs.microsoft.com/en-us/cpp/cpp/calling-conventions).
285 // For information on calling conventions for x64 processors, see
286 // Calling Convention
287 // (https://docs.microsoft.com/en-us/cpp/build/x64-calling-convention).
288 #pragma message("Alignment of class vec is not in accordance with SYCL \
289 specification requirements, a limitation of the MSVC compiler(Error C2719).\
290 Requested alignment applied, limited at 64.")
291 #define __SYCL_ALIGNED_VAR(type, x, var) \
292  type __declspec(align((x < 64) ? x : 64)) var
293 #else
294 #define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
295 #endif
296 
297 #endif
298 
299 template <typename Type, int NumElements> class vec {
304  using DataT = Type;
305 
306  // This represent type of underlying value. There should be only one field
307  // in the class, so vec<float, 16> should be equal to float16 in memory.
308  using DataType = typename detail::VecStorage<DataT, NumElements>::DataType;
309 
310  static constexpr bool IsHostHalf =
311  std::is_same_v<DataT, sycl::detail::half_impl::half> &&
312  std::is_same_v<sycl::detail::half_impl::StorageT,
313  sycl::detail::host_half_impl::half>;
314 
315 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
316 
317  static constexpr size_t AdjustedNum = (NumElements == 3) ? 4 : NumElements;
318  static constexpr size_t Sz = sizeof(DataT) * AdjustedNum;
319  static constexpr bool IsSizeGreaterThanMaxAlign =
320  (Sz > detail::MaxVecAlignment);
321 
322  // TODO: There is no support for vector half type on host yet.
323  // Also, when Sz is greater than alignment, we use std::array instead of
324  // vector extension. This is for MSVC compatibility, which has a max alignment
325  // of 64 for direct params. If we drop MSVC, we can have alignment the same as
326  // size and use vector extensions for all sizes.
327  static constexpr bool IsUsingArrayOnDevice =
328  (IsHostHalf || IsSizeGreaterThanMaxAlign);
329 
330 #if defined(__SYCL_DEVICE_ONLY__)
331  static constexpr bool NativeVec = NumElements > 1 && !IsUsingArrayOnDevice;
332  static constexpr bool IsUsingArrayOnHost = false; // not compiling for host.
333 #else
334  static constexpr bool NativeVec = false;
335  static constexpr bool IsUsingArrayOnHost = true; // host always std::array.
336 #endif
337 
338 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
339 
340 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
341  static constexpr bool IsUsingArrayOnDevice = IsHostHalf;
342 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
343 
344  static constexpr int getNumElements() { return NumElements; }
345 
346  // SizeChecker is needed for vec(const argTN &... args) ctor to validate args.
347  template <int Counter, int MaxValue, class...>
348  struct SizeChecker : std::conditional_t<Counter == MaxValue, std::true_type,
349  std::false_type> {};
350 
351  template <int Counter, int MaxValue, typename DataT_, class... tail>
352  struct SizeChecker<Counter, MaxValue, DataT_, tail...>
353  : std::conditional_t<Counter + 1 <= MaxValue,
354  SizeChecker<Counter + 1, MaxValue, tail...>,
355  std::false_type> {};
356 
357  // Utility trait for creating an std::array from an vector argument.
358  template <typename DataT_, typename T, std::size_t... Is>
359  static constexpr std::array<DataT_, sizeof...(Is)>
360  VecToArray(const vec<T, sizeof...(Is)> &V, std::index_sequence<Is...>) {
361  return {static_cast<DataT_>(V.getValue(Is))...};
362  }
363  template <typename DataT_, typename T, int N, typename T2, typename T3,
364  template <typename> class T4, int... T5, std::size_t... Is>
365  static constexpr std::array<DataT_, sizeof...(Is)>
366  VecToArray(const detail::SwizzleOp<vec<T, N>, T2, T3, T4, T5...> &V,
367  std::index_sequence<Is...>) {
368  return {static_cast<DataT_>(V.getValue(Is))...};
369  }
370  template <typename DataT_, typename T, int N, typename T2, typename T3,
371  template <typename> class T4, int... T5, std::size_t... Is>
372  static constexpr std::array<DataT_, sizeof...(Is)>
373  VecToArray(const detail::SwizzleOp<const vec<T, N>, T2, T3, T4, T5...> &V,
374  std::index_sequence<Is...>) {
375  return {static_cast<DataT_>(V.getValue(Is))...};
376  }
377  template <typename DataT_, typename T, int N>
378  static constexpr std::array<DataT_, N>
379  FlattenVecArgHelper(const vec<T, N> &A) {
380  return VecToArray<DataT_>(A, std::make_index_sequence<N>());
381  }
382  template <typename DataT_, typename T, int N, typename T2, typename T3,
383  template <typename> class T4, int... T5>
384  static constexpr std::array<DataT_, sizeof...(T5)> FlattenVecArgHelper(
385  const detail::SwizzleOp<vec<T, N>, T2, T3, T4, T5...> &A) {
386  return VecToArray<DataT_>(A, std::make_index_sequence<sizeof...(T5)>());
387  }
388  template <typename DataT_, typename T, int N, typename T2, typename T3,
389  template <typename> class T4, int... T5>
390  static constexpr std::array<DataT_, sizeof...(T5)> FlattenVecArgHelper(
391  const detail::SwizzleOp<const vec<T, N>, T2, T3, T4, T5...> &A) {
392  return VecToArray<DataT_>(A, std::make_index_sequence<sizeof...(T5)>());
393  }
394  template <typename DataT_, typename T>
395  static constexpr auto FlattenVecArgHelper(const T &A) {
396 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
397  return std::array<DataT_, 1>{vec_data<DataT_>::get(static_cast<DataT_>(A))};
398 #else
399  return std::array<DataT_, 1>{vec_data<DataT_>::get(A)};
400 #endif
401  }
402  template <typename DataT_, typename T> struct FlattenVecArg {
403  constexpr auto operator()(const T &A) const {
404  return FlattenVecArgHelper<DataT_>(A);
405  }
406  };
407 
408  // Alias for shortening the vec arguments to array converter.
409  template <typename DataT_, typename... ArgTN>
410  using VecArgArrayCreator =
411  detail::ArrayCreator<DataT_, FlattenVecArg, ArgTN...>;
412 
413 #define __SYCL_ALLOW_VECTOR_SIZES(num_elements) \
414  template <int Counter, int MaxValue, typename DataT_, class... tail> \
415  struct SizeChecker<Counter, MaxValue, vec<DataT_, num_elements>, tail...> \
416  : std::conditional_t< \
417  Counter + (num_elements) <= MaxValue, \
418  SizeChecker<Counter + (num_elements), MaxValue, tail...>, \
419  std::false_type> {}; \
420  template <int Counter, int MaxValue, typename DataT_, typename T2, \
421  typename T3, template <typename> class T4, int... T5, \
422  class... tail> \
423  struct SizeChecker< \
424  Counter, MaxValue, \
425  detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
426  tail...> \
427  : std::conditional_t< \
428  Counter + sizeof...(T5) <= MaxValue, \
429  SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
430  std::false_type> {}; \
431  template <int Counter, int MaxValue, typename DataT_, typename T2, \
432  typename T3, template <typename> class T4, int... T5, \
433  class... tail> \
434  struct SizeChecker< \
435  Counter, MaxValue, \
436  detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
437  tail...> \
438  : std::conditional_t< \
439  Counter + sizeof...(T5) <= MaxValue, \
440  SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
441  std::false_type> {};
442 
443  __SYCL_ALLOW_VECTOR_SIZES(1)
444  __SYCL_ALLOW_VECTOR_SIZES(2)
445  __SYCL_ALLOW_VECTOR_SIZES(3)
446  __SYCL_ALLOW_VECTOR_SIZES(4)
447  __SYCL_ALLOW_VECTOR_SIZES(8)
448  __SYCL_ALLOW_VECTOR_SIZES(16)
449 #undef __SYCL_ALLOW_VECTOR_SIZES
450 
451  // TypeChecker is needed for vec(const argTN &... args) ctor to validate args.
452  template <typename T, typename DataT_>
453  struct TypeChecker : std::is_convertible<T, DataT_> {};
454 #define __SYCL_ALLOW_VECTOR_TYPES(num_elements) \
455  template <typename DataT_> \
456  struct TypeChecker<vec<DataT_, num_elements>, DataT_> : std::true_type {}; \
457  template <typename DataT_, typename T2, typename T3, \
458  template <typename> class T4, int... T5> \
459  struct TypeChecker< \
460  detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, DataT_> \
461  : std::true_type {}; \
462  template <typename DataT_, typename T2, typename T3, \
463  template <typename> class T4, int... T5> \
464  struct TypeChecker< \
465  detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
466  DataT_> : std::true_type {};
467 
468  __SYCL_ALLOW_VECTOR_TYPES(1)
469  __SYCL_ALLOW_VECTOR_TYPES(2)
470  __SYCL_ALLOW_VECTOR_TYPES(3)
471  __SYCL_ALLOW_VECTOR_TYPES(4)
472  __SYCL_ALLOW_VECTOR_TYPES(8)
473  __SYCL_ALLOW_VECTOR_TYPES(16)
474 #undef __SYCL_ALLOW_VECTOR_TYPES
475 
476  template <int... Indexes>
477  using Swizzle =
478  detail::SwizzleOp<vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
479  detail::GetOp, Indexes...>;
480 
481  template <int... Indexes>
482  using ConstSwizzle =
483  detail::SwizzleOp<const vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
484  detail::GetOp, Indexes...>;
485 
486  // Shortcuts for args validation in vec(const argTN &... args) ctor.
487  template <typename... argTN>
488  using EnableIfSuitableTypes = typename std::enable_if_t<
489  std::conjunction_v<TypeChecker<argTN, DataT>...>>;
490 
491  template <typename... argTN>
492  using EnableIfSuitableNumElements =
493  typename std::enable_if_t<SizeChecker<0, NumElements, argTN...>::value>;
494 
495  template <size_t... Is>
496  constexpr vec(const std::array<vec_data_t<DataT>, NumElements> &Arr,
497  std::index_sequence<Is...>)
498  : m_Data{vec_data_t<DataT>(static_cast<DataT>(Arr[Is]))...} {}
499 
500 public:
501  using element_type = DataT;
502  using rel_t = detail::rel_t<DataT>;
503 
504 #ifdef __SYCL_DEVICE_ONLY__
505 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
506  using vector_t =
507  typename detail::VecStorage<DataT, NumElements>::VectorDataType;
508 #else
509  using vector_t = DataType;
510 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
511 #endif // __SYCL_DEVICE_ONLY__
512 
513  vec() = default;
514 
515 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
516  constexpr vec(const vec &Rhs) = default;
517  constexpr vec(vec &&Rhs) = default;
518 
519  constexpr vec &operator=(const vec &Rhs) = default;
520 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
521 
522 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
523  // TODO Remove this difference between host and device side after
524  // when root cause of API incompatibility will be fixed
525 #ifdef __SYCL_DEVICE_ONLY__
526  vec(const vec &Rhs) = default;
527  vec &operator=(const vec &Rhs) = default;
528 #else
529  constexpr vec(const vec &Rhs) : m_Data(Rhs.m_Data) {}
530  constexpr vec &operator=(const vec &Rhs) = default;
531 #endif
532 
533  vec(vec &&Rhs) = default;
534 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
535 
536  // W/o this, things like "vec<char,*> = vec<signed char, *>" doesn't work.
537  template <typename Ty = DataT>
538  typename std::enable_if_t<!std::is_same_v<Ty, rel_t> &&
539  std::is_convertible_v<vec_data_t<Ty>, rel_t>,
540  vec &>
541  operator=(const vec<rel_t, NumElements> &Rhs) {
542  *this = Rhs.template as<vec>();
543  return *this;
544  }
545 
546 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
547  template <typename T = void>
548  using EnableIfNotHostHalf = typename std::enable_if_t<!IsHostHalf, T>;
549 
550  template <typename T = void>
551  using EnableIfHostHalf = typename std::enable_if_t<IsHostHalf, T>;
552 
553  template <typename T = void>
554  using EnableIfUsingArrayOnDevice =
555  typename std::enable_if_t<IsUsingArrayOnDevice, T>;
556 
557  template <typename T = void>
558  using EnableIfNotUsingArrayOnDevice =
559  typename std::enable_if_t<!IsUsingArrayOnDevice, T>;
560 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
561 
562 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
563 
564  template <typename T = void>
565  using EnableIfUsingArray =
566  typename std::enable_if_t<IsUsingArrayOnDevice || IsUsingArrayOnHost, T>;
567 
568  template <typename T = void>
569  using EnableIfNotUsingArray =
570  typename std::enable_if_t<!IsUsingArrayOnDevice && !IsUsingArrayOnHost,
571  T>;
572 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
573 
574 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
575 
576  template <typename Ty = DataT>
577  explicit constexpr vec(const EnableIfNotUsingArrayOnDevice<Ty> &arg)
578  : m_Data{DataType(vec_data<Ty>::get(arg))} {}
579 
580  template <typename Ty = DataT>
581  typename std::enable_if_t<
582  std::is_fundamental_v<vec_data_t<Ty>> ||
583  std::is_same_v<typename std::remove_const_t<Ty>, half>,
584  vec &>
585  operator=(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) {
586  m_Data = (DataType)vec_data<Ty>::get(Rhs);
587  return *this;
588  }
589 
590  template <typename Ty = DataT>
591  explicit constexpr vec(const EnableIfUsingArrayOnDevice<Ty> &arg)
592  : vec{detail::RepeatValue<NumElements>(
593  static_cast<vec_data_t<DataT>>(arg)),
594  std::make_index_sequence<NumElements>()} {}
595 
596  template <typename Ty = DataT>
597  typename std::enable_if_t<
598  std::is_fundamental_v<vec_data_t<Ty>> ||
599  std::is_same_v<typename std::remove_const_t<Ty>, half>,
600  vec &>
601  operator=(const EnableIfUsingArrayOnDevice<Ty> &Rhs) {
602  for (int i = 0; i < NumElements; ++i) {
603  setValue(i, Rhs);
604  }
605  return *this;
606  }
607 #else // __SYCL_USE_EXT_VECTOR_TYPE__
608  explicit constexpr vec(const DataT &arg)
609  : vec{detail::RepeatValue<NumElements>(
610  static_cast<vec_data_t<DataT>>(arg)),
611  std::make_index_sequence<NumElements>()} {}
612 
613  template <typename Ty = DataT>
614  typename std::enable_if_t<
615  std::is_fundamental_v<vec_data_t<Ty>> ||
616  std::is_same_v<typename std::remove_const_t<Ty>, half>,
617  vec &>
618  operator=(const DataT &Rhs) {
619  for (int i = 0; i < NumElements; ++i) {
620  setValue(i, Rhs);
621  }
622  return *this;
623  }
624 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
625 
626 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
627  // Optimized naive constructors with NumElements of DataT values.
628  // We don't expect compilers to optimize vararg recursive functions well.
629 
630  // Helper type to make specific constructors available only for specific
631  // number of elements.
632  template <int IdxNum, typename T = void>
633  using EnableIfMultipleElems = typename std::enable_if_t<
634  std::is_convertible_v<T, DataT> && NumElements == IdxNum, DataT>;
635  template <typename Ty = DataT>
636  constexpr vec(const EnableIfMultipleElems<2, Ty> Arg0,
637  const EnableIfNotUsingArrayOnDevice<Ty> Arg1)
638  : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1)} {}
639  template <typename Ty = DataT>
640  constexpr vec(const EnableIfMultipleElems<3, Ty> Arg0,
641  const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2)
642  : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
643  vec_data<Ty>::get(Arg2)} {}
644  template <typename Ty = DataT>
645  constexpr vec(const EnableIfMultipleElems<4, Ty> Arg0,
646  const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2,
647  const Ty Arg3)
648  : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
649  vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3)} {}
650  template <typename Ty = DataT>
651  constexpr vec(const EnableIfMultipleElems<8, Ty> Arg0,
652  const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2,
653  const DataT Arg3, const DataT Arg4, const DataT Arg5,
654  const DataT Arg6, const DataT Arg7)
655  : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
656  vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3),
657  vec_data<Ty>::get(Arg4), vec_data<Ty>::get(Arg5),
658  vec_data<Ty>::get(Arg6), vec_data<Ty>::get(Arg7)} {}
659  template <typename Ty = DataT>
660  constexpr vec(const EnableIfMultipleElems<16, Ty> Arg0,
661  const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2,
662  const DataT Arg3, const DataT Arg4, const DataT Arg5,
663  const DataT Arg6, const DataT Arg7, const DataT Arg8,
664  const DataT Arg9, const DataT ArgA, const DataT ArgB,
665  const DataT ArgC, const DataT ArgD, const DataT ArgE,
666  const DataT ArgF)
667  : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
668  vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3),
669  vec_data<Ty>::get(Arg4), vec_data<Ty>::get(Arg5),
670  vec_data<Ty>::get(Arg6), vec_data<Ty>::get(Arg7),
671  vec_data<Ty>::get(Arg8), vec_data<Ty>::get(Arg9),
672  vec_data<Ty>::get(ArgA), vec_data<Ty>::get(ArgB),
673  vec_data<Ty>::get(ArgC), vec_data<Ty>::get(ArgD),
674  vec_data<Ty>::get(ArgE), vec_data<Ty>::get(ArgF)} {}
675 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
676 
677  // Constructor from values of base type or vec of base type. Checks that
678  // base types are match and that the NumElements == sum of lengths of args.
679  template <typename... argTN, typename = EnableIfSuitableTypes<argTN...>,
680  typename = EnableIfSuitableNumElements<argTN...>>
681  constexpr vec(const argTN &...args)
682  : vec{VecArgArrayCreator<vec_data_t<DataT>, argTN...>::Create(args...),
683  std::make_index_sequence<NumElements>()} {}
684 
685 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
686 
687 #ifdef __SYCL_DEVICE_ONLY__
688  template <typename vector_t_ = vector_t,
689  typename =
690  typename std::enable_if_t<std::is_same_v<vector_t_, vector_t> &&
691  !std::is_same_v<vector_t_, DataT>>>
692  constexpr vec(vector_t openclVector) {
693  if constexpr (!IsUsingArrayOnDevice) {
694  m_Data = openclVector;
695  } else {
696  m_Data = bit_cast<DataType>(openclVector);
697  }
698  }
699 
700  operator vector_t() const {
701  if constexpr (!IsUsingArrayOnDevice) {
702  return m_Data;
703  } else {
704  auto ptr = bit_cast<const vector_t *>((&m_Data)->data());
705  return *ptr;
706  }
707  }
708 #endif // __SYCL_DEVICE_ONLY__
709 
710 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
711 
712 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
713 
714 #ifdef __SYCL_DEVICE_ONLY__
715  template <typename vector_t_ = vector_t,
716  typename =
717  typename std::enable_if_t<std::is_same_v<vector_t_, vector_t> &&
718  !std::is_same_v<vector_t_, DataT>>>
719  constexpr vec(vector_t openclVector) : m_Data(openclVector) {}
720  operator vector_t() const { return m_Data; }
721 #endif
722 
723 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
724 
725  // Available only when: NumElements == 1
726  template <int N = NumElements>
727  operator typename std::enable_if_t<N == 1, DataT>() const {
728  return vec_data<DataT>::get(m_Data);
729  }
730 
731  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
732  static constexpr size_t get_count() { return size(); }
733  static constexpr size_t size() noexcept { return NumElements; }
734  __SYCL2020_DEPRECATED(
735  "get_size() is deprecated, please use byte_size() instead")
736  static constexpr size_t get_size() { return byte_size(); }
737  static constexpr size_t byte_size() noexcept { return sizeof(m_Data); }
738 
739  // convertImpl can't be called with the same From and To types and therefore
740  // we need this version of convert which is mostly no-op.
741  template <typename convertT,
742  rounding_mode roundingMode = rounding_mode::automatic>
743  std::enable_if_t<
744  std::is_same_v<vec_data_t<DataT>, vec_data_t<convertT>> ||
745  std::is_same_v<detail::ConvertToOpenCLType_t<vec_data_t<DataT>>,
746  detail::ConvertToOpenCLType_t<vec_data_t<convertT>>>,
747  vec<convertT, NumElements>>
748  convert() const {
749  static_assert(std::is_integral_v<vec_data_t<convertT>> ||
750  detail::is_floating_point<convertT>::value,
751  "Unsupported convertT");
752  if constexpr (!std::is_same_v<DataT, convertT>) {
753  // Dummy conversion for cases like vec<signed char> -> vec<char>
754  vec<convertT, NumElements> Result;
755  for (size_t I = 0; I < NumElements; ++I)
756  Result.setValue(I, static_cast<convertT>(getValue(I)));
757 
758  return Result;
759  } else {
760  // No conversion necessary
761  return *this;
762  }
763  }
764 
765  template <typename convertT,
766  rounding_mode roundingMode = rounding_mode::automatic>
767  std::enable_if_t<
768  !std::is_same_v<vec_data_t<DataT>, vec_data_t<convertT>> &&
769  !std::is_same_v<detail::ConvertToOpenCLType_t<vec_data_t<DataT>>,
770  detail::ConvertToOpenCLType_t<vec_data_t<convertT>>>,
771  vec<convertT, NumElements>>
772  convert() const {
773  static_assert(std::is_integral_v<vec_data_t<convertT>> ||
774  detail::is_floating_point<convertT>::value,
775  "Unsupported convertT");
776  using T = vec_data_t<DataT>;
777  using R = vec_data_t<convertT>;
778  using OpenCLT = detail::ConvertToOpenCLType_t<T>;
779  using OpenCLR = detail::ConvertToOpenCLType_t<R>;
780  vec<convertT, NumElements> Result;
781 
782 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
783  using OpenCLVecT = OpenCLT __attribute__((ext_vector_type(NumElements)));
784  using OpenCLVecR = OpenCLR __attribute__((ext_vector_type(NumElements)));
785  // Whole vector conversion can only be done, if:
786  constexpr bool canUseNativeVectorConvert =
787 #ifdef __NVPTX__
788  // - we are not on CUDA, see intel/llvm#11840
789  false &&
790 #endif
791  // - both vectors are represented using native vector types;
792  NativeVec && vec<convertT, NumElements>::NativeVec &&
793  // - vec storage has an equivalent OpenCL native vector it is implicitly
794  // convertible to. There are some corner cases where it is not the
795  // case with char, long and long long types.
796  std::is_convertible_v<decltype(m_Data), OpenCLVecT> &&
797  std::is_convertible_v<decltype(Result.m_Data), OpenCLVecR> &&
798  // - it is not a signed to unsigned (or vice versa) conversion
799  // see comments within 'convertImpl' for more details;
800  !detail::is_sint_to_from_uint<T, R>::value &&
801  // - destination type is not bool. bool is stored as integer under the
802  // hood and therefore conversion to bool looks like conversion between
803  // two integer types. Since bit pattern for true and false is not
804  // defined, there is no guarantee that integer conversion yields
805  // right results here;
806  !std::is_same_v<convertT, bool>;
807  if constexpr (canUseNativeVectorConvert) {
808  Result.m_Data = detail::convertImpl<T, R, roundingMode, NumElements,
809  OpenCLVecT, OpenCLVecR>(m_Data);
810  } else
811 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) &&
812  // defined(__SYCL_DEVICE_ONLY__)
813  {
814  // Otherwise, we fallback to per-element conversion:
815  for (size_t I = 0; I < NumElements; ++I) {
816  Result.setValue(
817  I, vec_data<convertT>::get(
818  detail::convertImpl<T, R, roundingMode, 1, OpenCLT, OpenCLR>(
819  vec_data<DataT>::get(getValue(I)))));
820  }
821  }
822 
823  return Result;
824  }
825 
826  template <typename asT> asT as() const {
827  static_assert((sizeof(*this) == sizeof(asT)),
828  "The new SYCL vec type must have the same storage size in "
829  "bytes as this SYCL vec");
830  static_assert(
831  detail::is_contained<asT, detail::gtl::vector_basic_list>::value ||
832  detail::is_contained<asT, detail::gtl::vector_bool_list>::value,
833  "asT must be SYCL vec of a different element type and "
834  "number of elements specified by asT");
835  asT Result;
836  detail::memcpy(&Result.m_Data, &m_Data, sizeof(decltype(Result.m_Data)));
837  return Result;
838  }
839 
840  template <int... SwizzleIndexes> Swizzle<SwizzleIndexes...> swizzle() {
841  return this;
842  }
843 
844  template <int... SwizzleIndexes>
845  ConstSwizzle<SwizzleIndexes...> swizzle() const {
846  return this;
847  }
848 
849  // ext_vector_type is used as an underlying type for sycl::vec on device.
850  // The problem is that for clang vector types the return of operator[] is a
851  // temporary and not a reference to the element in the vector. In practice
852  // reinterpret_cast<DataT *>(&m_Data)[i]; is working. According to
853  // http://llvm.org/docs/GetElementPtr.html#can-gep-index-into-vector-elements
854  // this is not disallowed now. But could probably be disallowed in the future.
855  // That is why tests are added to check that behavior of the compiler has
856  // not changed.
857  //
858  // Implement operator [] in the same way for host and device.
859  // TODO: change host side implementation when underlying type for host side
860  // will be changed to std::array.
861  const DataT &operator[](int i) const {
862  return reinterpret_cast<const DataT *>(&m_Data)[i];
863  }
864 
865  DataT &operator[](int i) { return reinterpret_cast<DataT *>(&m_Data)[i]; }
866 
867  // Begin hi/lo, even/odd, xyzw, and rgba swizzles.
868 private:
869  // Indexer used in the swizzles.def
870  // Currently it is defined as a template struct. Replacing it with a constexpr
871  // function would activate a bug in MSVC that is fixed only in v19.20.
872  // Until then MSVC does not recognize such constexpr functions as const and
873  // thus does not let using them in template parameters inside swizzle.def.
874  template <int Index> struct Indexer {
875  static constexpr int value = Index;
876  };
877 
878 public:
879 #ifdef __SYCL_ACCESS_RETURN
880 #error "Undefine __SYCL_ACCESS_RETURN macro"
881 #endif
882 #define __SYCL_ACCESS_RETURN this
883 #include "swizzles.def"
884 #undef __SYCL_ACCESS_RETURN
885  // End of hi/lo, even/odd, xyzw, and rgba swizzles.
886 
887  template <access::address_space Space, access::decorated DecorateAddress>
888  void load(size_t Offset, multi_ptr<const DataT, Space, DecorateAddress> Ptr) {
889  for (int I = 0; I < NumElements; I++) {
890  setValue(I, *multi_ptr<const DataT, Space, DecorateAddress>(
891  Ptr + Offset * NumElements + I));
892  }
893  }
894  template <access::address_space Space, access::decorated DecorateAddress>
895  void load(size_t Offset, multi_ptr<DataT, Space, DecorateAddress> Ptr) {
896  multi_ptr<const DataT, Space, DecorateAddress> ConstPtr(Ptr);
897  load(Offset, ConstPtr);
898  }
899  template <int Dimensions, access::mode Mode,
900  access::placeholder IsPlaceholder, access::target Target,
901  typename PropertyListT>
902  void
903  load(size_t Offset,
904  accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
905  Acc) {
906  multi_ptr<const DataT, detail::TargetToAS<Target>::AS,
907  access::decorated::yes>
908  MultiPtr(Acc);
909  load(Offset, MultiPtr);
910  }
911  template <access::address_space Space, access::decorated DecorateAddress>
912  void store(size_t Offset,
913  multi_ptr<DataT, Space, DecorateAddress> Ptr) const {
914  for (int I = 0; I < NumElements; I++) {
915  *multi_ptr<DataT, Space, DecorateAddress>(Ptr + Offset * NumElements +
916  I) = getValue(I);
917  }
918  }
919  template <int Dimensions, access::mode Mode,
920  access::placeholder IsPlaceholder, access::target Target,
921  typename PropertyListT>
922  void
923  store(size_t Offset,
924  accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
925  Acc) {
926  multi_ptr<DataT, detail::TargetToAS<Target>::AS, access::decorated::yes>
927  MultiPtr(Acc);
928  store(Offset, MultiPtr);
929  }
930 
931  void ConvertToDataT() {
932  for (size_t i = 0; i < NumElements; ++i) {
933  DataT tmp = getValue(i);
934  setValue(i, tmp);
935  }
936  }
937 
938 #ifdef __SYCL_BINOP
939 #error "Undefine __SYCL_BINOP macro"
940 #endif
941 
942 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
943 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \
944  template <typename Ty = vec> \
945  vec operator BINOP(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) const { \
946  vec Ret; \
947  Ret.m_Data = m_Data BINOP Rhs.m_Data; \
948  if constexpr (std::is_same_v<Type, bool> && CONVERT) { \
949  Ret.ConvertToDataT(); \
950  } \
951  return Ret; \
952  } \
953  template <typename Ty = vec> \
954  vec operator BINOP(const EnableIfUsingArrayOnDevice<Ty> &Rhs) const { \
955  vec Ret; \
956  for (size_t I = 0; I < NumElements; ++I) { \
957  Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
958  } \
959  return Ret; \
960  } \
961  template <typename T> \
962  typename std::enable_if_t< \
963  std::is_convertible_v<DataT, T> && \
964  (std::is_fundamental_v<vec_data_t<T>> || \
965  std::is_same_v<typename std::remove_const_t<T>, half>), \
966  vec> \
967  operator BINOP(const T & Rhs) const { \
968  return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
969  } \
970  vec &operator OPASSIGN(const vec & Rhs) { \
971  *this = *this BINOP Rhs; \
972  return *this; \
973  } \
974  template <int Num = NumElements> \
975  typename std::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
976  const DataT & Rhs) { \
977  *this = *this BINOP vec(Rhs); \
978  return *this; \
979  }
980 #else // __SYCL_USE_EXT_VECTOR_TYPE__
981 
982 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
983 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \
984  vec operator BINOP(const vec &Rhs) const { \
985  vec Ret{}; \
986  if constexpr (NativeVec) \
987  Ret.m_Data = m_Data BINOP Rhs.m_Data; \
988  else \
989  for (size_t I = 0; I < NumElements; ++I) \
990  Ret.setValue(I, (DataT)(vec_data<DataT>::get(getValue( \
991  I)) BINOP vec_data<DataT>::get(Rhs.getValue(I)))); \
992  return Ret; \
993  } \
994  template <typename T> \
995  typename std::enable_if_t< \
996  std::is_convertible_v<DataT, T> && \
997  (std::is_fundamental_v<vec_data_t<T>> || \
998  std::is_same_v<typename std::remove_const_t<T>, half>), \
999  vec> \
1000  operator BINOP(const T & Rhs) const { \
1001  return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1002  } \
1003  vec &operator OPASSIGN(const vec & Rhs) { \
1004  *this = *this BINOP Rhs; \
1005  return *this; \
1006  } \
1007  template <int Num = NumElements> \
1008  typename std::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1009  const DataT & Rhs) { \
1010  *this = *this BINOP vec(Rhs); \
1011  return *this; \
1012  }
1013 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1014 
1015 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1016 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \
1017  vec operator BINOP(const vec &Rhs) const { \
1018  vec Ret; \
1019  for (size_t I = 0; I < NumElements; ++I) { \
1020  Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
1021  } \
1022  return Ret; \
1023  } \
1024  template <typename T> \
1025  typename std::enable_if_t< \
1026  std::is_convertible_v<DataT, T> && \
1027  (std::is_fundamental_v<vec_data_t<T>> || \
1028  std::is_same_v<typename std::remove_const_t<T>, half>), \
1029  vec> \
1030  operator BINOP(const T & Rhs) const { \
1031  return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1032  } \
1033  vec &operator OPASSIGN(const vec & Rhs) { \
1034  *this = *this BINOP Rhs; \
1035  return *this; \
1036  } \
1037  template <int Num = NumElements> \
1038  typename std::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1039  const DataT & Rhs) { \
1040  *this = *this BINOP vec(Rhs); \
1041  return *this; \
1042  }
1043 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1044 
1045 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
1046 
1047  __SYCL_BINOP(+, +=, true)
1048  __SYCL_BINOP(-, -=, true)
1049  __SYCL_BINOP(*, *=, false)
1050  __SYCL_BINOP(/, /=, false)
1051 
1052  // TODO: The following OPs are available only when: DataT != cl_float &&
1053  // DataT != cl_double && DataT != cl_half
1054  __SYCL_BINOP(%, %=, false)
1055  __SYCL_BINOP(|, |=, false)
1056  __SYCL_BINOP(&, &=, false)
1057  __SYCL_BINOP(^, ^=, false)
1058  __SYCL_BINOP(>>, >>=, false)
1059  __SYCL_BINOP(<<, <<=, true)
1060 #undef __SYCL_BINOP
1061 #undef __SYCL_BINOP_HELP
1062 
1063  // Note: vec<>/SwizzleOp logical value is 0/-1 logic, as opposed to 0/1 logic.
1064  // As far as CTS validation is concerned, 0/-1 logic also applies when
1065  // NumElements is equal to one, which is somewhat inconsistent with being
1066  // transparent with scalar data.
1067  // TODO: Determine if vec<, NumElements=1> is needed at all, remove this
1068  // inconsistency if not by disallowing one-element vectors (as in OpenCL)
1069 
1070 #ifdef __SYCL_RELLOGOP
1071 #error "Undefine __SYCL_RELLOGOP macro"
1072 #endif
1073 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1074 // by SYCL device compiler only.
1075 #ifdef __SYCL_DEVICE_ONLY__
1076 #define __SYCL_RELLOGOP(RELLOGOP) \
1077  vec<rel_t, NumElements> operator RELLOGOP(const vec & Rhs) const { \
1078  vec<rel_t, NumElements> Ret{}; \
1079  /* This special case is needed since there are no standard operator|| */ \
1080  /* or operator&& functions for std::array. */ \
1081  if constexpr (IsUsingArrayOnDevice && \
1082  (std::string_view(#RELLOGOP) == "||" || \
1083  std::string_view(#RELLOGOP) == "&&")) { \
1084  for (size_t I = 0; I < NumElements; ++I) { \
1085  Ret.setValue(I, \
1086  -(vec_data<DataT>::get(getValue(I)) \
1087  RELLOGOP vec_data<DataT>::get(Rhs.getValue(I)))); \
1088  } \
1089  } else { \
1090  Ret = vec<rel_t, NumElements>( \
1091  (typename vec<rel_t, NumElements>::vector_t)( \
1092  m_Data RELLOGOP Rhs.m_Data)); \
1093  if (NumElements == 1) /*Scalar 0/1 logic was applied, invert*/ \
1094  Ret *= -1; \
1095  } \
1096  return Ret; \
1097  } \
1098  template <typename T> \
1099  typename std::enable_if_t<std::is_convertible_v<T, DataT> && \
1100  (std::is_fundamental_v<vec_data_t<T>> || \
1101  std::is_same_v<T, half>), \
1102  vec<rel_t, NumElements>> \
1103  operator RELLOGOP(const T & Rhs) const { \
1104  return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1105  }
1106 #else
1107 #define __SYCL_RELLOGOP(RELLOGOP) \
1108  vec<rel_t, NumElements> operator RELLOGOP(const vec & Rhs) const { \
1109  vec<rel_t, NumElements> Ret{}; \
1110  for (size_t I = 0; I < NumElements; ++I) { \
1111  Ret.setValue(I, -(vec_data<DataT>::get(getValue(I)) \
1112  RELLOGOP vec_data<DataT>::get(Rhs.getValue(I)))); \
1113  } \
1114  return Ret; \
1115  } \
1116  template <typename T> \
1117  typename std::enable_if_t<std::is_convertible_v<T, DataT> && \
1118  (std::is_fundamental_v<vec_data_t<T>> || \
1119  std::is_same_v<T, half>), \
1120  vec<rel_t, NumElements>> \
1121  operator RELLOGOP(const T & Rhs) const { \
1122  return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1123  }
1124 #endif
1125 
1126  __SYCL_RELLOGOP(==)
1127  __SYCL_RELLOGOP(!=)
1128  __SYCL_RELLOGOP(>)
1129  __SYCL_RELLOGOP(<)
1130  __SYCL_RELLOGOP(>=)
1131  __SYCL_RELLOGOP(<=)
1132  // TODO: limit to integral types.
1133  __SYCL_RELLOGOP(&&)
1134  __SYCL_RELLOGOP(||)
1135 #undef __SYCL_RELLOGOP
1136 
1137 #ifdef __SYCL_UOP
1138 #error "Undefine __SYCL_UOP macro"
1139 #endif
1140 #define __SYCL_UOP(UOP, OPASSIGN) \
1141  vec &operator UOP() { \
1142  *this OPASSIGN vec_data<DataT>::get(1); \
1143  return *this; \
1144  } \
1145  vec operator UOP(int) { \
1146  vec Ret(*this); \
1147  *this OPASSIGN vec_data<DataT>::get(1); \
1148  return Ret; \
1149  }
1150 
1151  __SYCL_UOP(++, +=)
1152  __SYCL_UOP(--, -=)
1153 #undef __SYCL_UOP
1154 
1155 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1156 
1157  // operator~() available only when: dataT != float && dataT != double
1158  // && dataT != half
1159  template <typename T = DataT>
1160  typename std::enable_if_t<!std::is_floating_point_v<vec_data_t<T>> &&
1161  (!IsUsingArrayOnDevice && !IsUsingArrayOnHost),
1162  vec>
1163  operator~() const {
1164  vec Ret{(typename vec::DataType) ~m_Data};
1165  if constexpr (std::is_same_v<Type, bool>) {
1166  Ret.ConvertToDataT();
1167  }
1168  return Ret;
1169  }
1170  template <typename T = DataT>
1171  typename std::enable_if_t<!std::is_floating_point_v<vec_data_t<T>> &&
1172  (IsUsingArrayOnDevice || IsUsingArrayOnHost),
1173  vec>
1174  operator~() const {
1175  vec Ret{};
1176  for (size_t I = 0; I < NumElements; ++I) {
1177  Ret.setValue(I, ~getValue(I));
1178  }
1179  return Ret;
1180  }
1181 
1182  template <typename T>
1183 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1184  using OpNotRet = detail::rel_t<T>;
1185 #else
1186  using OpNotRet = T;
1187 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
1188 
1189  // operator!
1190  template <typename T = DataT, int N = NumElements>
1191  EnableIfNotUsingArray<vec<OpNotRet<T>, N>> operator!() const {
1192  return vec<T, N>{(typename vec<DataT, NumElements>::DataType) !m_Data}
1193 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1194  .template as<vec<OpNotRet<T>, N>>();
1195 #else
1196  ;
1197 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
1198  }
1199 
1200  // std::byte neither supports ! unary op or casting, so special handling is
1201  // needed. And, worse, Windows has a conflict with 'byte'.
1202 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1203  template <typename T = DataT, int N = NumElements>
1204  typename std::enable_if_t<std::is_same_v<std::byte, T> &&
1205  (IsUsingArrayOnDevice || IsUsingArrayOnHost),
1206  vec<OpNotRet<T>, N>>
1207  operator!() const {
1208  vec Ret{};
1209  for (size_t I = 0; I < NumElements; ++I) {
1210  Ret.setValue(I, std::byte{!vec_data<DataT>::get(getValue(I))});
1211  }
1212 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1213  return Ret.template as<vec<OpNotRet<T>, N>>();
1214 #else
1215  return Ret;
1216 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
1217  }
1218 
1219  template <typename T = DataT, int N = NumElements>
1220  typename std::enable_if_t<!std::is_same_v<std::byte, T> &&
1221  (IsUsingArrayOnDevice || IsUsingArrayOnHost),
1222  vec<OpNotRet<T>, N>>
1223  operator!() const {
1224  vec Ret{};
1225  for (size_t I = 0; I < NumElements; ++I)
1226  Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1227 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1228  return Ret.template as<vec<OpNotRet<T>, N>>();
1229 #else
1230  return Ret;
1231 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
1232  }
1233 #else
1234  template <typename T = DataT, int N = NumElements>
1235  EnableIfUsingArray<vec<OpNotRet<T>, N>> operator!() const {
1236  vec Ret{};
1237  for (size_t I = 0; I < NumElements; ++I)
1238  Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1239 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1240  return Ret.template as<vec<OpNotRet<T>, N>>();
1241 #else
1242  return Ret;
1243 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
1244  }
1245 #endif
1246 
1247  // operator +
1248  template <typename T = vec> EnableIfNotUsingArray<T> operator+() const {
1249  return vec{+m_Data};
1250  }
1251 
1252  template <typename T = vec> EnableIfUsingArray<T> operator+() const {
1253  vec Ret{};
1254  for (size_t I = 0; I < NumElements; ++I)
1255  Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1256  return Ret;
1257  }
1258 
1259  // operator -
1260  template <typename T = vec> EnableIfNotUsingArray<T> operator-() const {
1261  vec Ret{-m_Data};
1262  if constexpr (std::is_same_v<Type, bool>) {
1263  Ret.ConvertToDataT();
1264  }
1265  return Ret;
1266  }
1267 
1268  template <typename T = vec> EnableIfUsingArray<T> operator-() const {
1269  vec Ret{};
1270  for (size_t I = 0; I < NumElements; ++I)
1271  Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
1272  return Ret;
1273  }
1274 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1275 
1276 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1277 
1278  // Available only when: dataT != cl_float && dataT != cl_double
1279  // && dataT != cl_half
1280  template <typename T = DataT>
1281  typename std::enable_if_t<std::is_integral_v<vec_data_t<T>>, vec>
1282  operator~() const {
1283 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1284 // by SYCL device compiler only.
1285 #ifdef __SYCL_DEVICE_ONLY__
1286  vec Ret{(typename vec::DataType) ~m_Data};
1287  if constexpr (std::is_same_v<Type, bool>) {
1288  Ret.ConvertToDataT();
1289  }
1290  return Ret;
1291 #else
1292  vec Ret{};
1293  for (size_t I = 0; I < NumElements; ++I) {
1294  Ret.setValue(I, ~getValue(I));
1295  }
1296  return Ret;
1297 #endif
1298  }
1299 
1300 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1301  template <typename T = DataT, int N = NumElements>
1302  typename std::enable_if_t<std::is_same_v<std::byte, T>, vec<T, N>>
1303  operator!() const {
1304 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1305 // by SYCL device compiler only.
1306 #ifdef __SYCL_DEVICE_ONLY__
1307  return vec<T, N>{(typename vec<T, N>::DataType) !m_Data};
1308 #else
1309  vec<T, N> Ret;
1310  for (size_t I = 0; I < N; ++I) {
1311  Ret.setValue(I, std::byte{!vec_data<DataT>::get(getValue(I))});
1312  }
1313  return Ret;
1314 #endif
1315  }
1316 
1317  template <typename T = DataT, int N = NumElements>
1318  typename std::enable_if_t<!std::is_same_v<std::byte, T>, vec<T, N>>
1319  operator!() const {
1320 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1321 // by SYCL device compiler only.
1322 #ifdef __SYCL_DEVICE_ONLY__
1323  return vec<T, N>{(typename vec<T, N>::DataType) !m_Data}; // T, N
1324 #else
1325  vec<T, N> Ret;
1326  for (size_t I = 0; I < N; ++I) {
1327  Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1328  }
1329  return Ret;
1330 #endif
1331  }
1332 #else // no std::byte
1333  template <typename T = DataT, int N = NumElements>
1334  vec<T, N> operator!() const {
1335 #ifdef __SYCL_DEVICE_ONLY__
1336  return vec<T, N>{(typename vec<T, N>::DataType) !m_Data}; // T, N
1337 #else
1338  vec<T, N> Ret;
1339  for (size_t I = 0; I < N; ++I) {
1340  Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1341  }
1342  return Ret;
1343 #endif
1344  }
1345 #endif // std::byte check
1346 
1347  vec operator+() const {
1348 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1349 // by SYCL device compiler only.
1350 #ifdef __SYCL_DEVICE_ONLY__
1351  return vec{+m_Data};
1352 #else
1353  vec Ret{};
1354  for (size_t I = 0; I < NumElements; ++I)
1355  Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1356  return Ret;
1357 #endif
1358  }
1359 
1360  vec operator-() const {
1361 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1362 // by SYCL device compiler only.
1363 #ifdef __SYCL_DEVICE_ONLY__
1364  vec Ret{-m_Data};
1365  if constexpr (std::is_same_v<Type, bool>) {
1366  Ret.ConvertToDataT();
1367  }
1368  return Ret;
1369 #else
1370  vec Ret{};
1371  for (size_t I = 0; I < NumElements; ++I)
1372  Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
1373  return Ret;
1374 #endif
1375  }
1376 
1377 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1378 
1379  // CP ---------------
1380 
1381  // OP is: &&, ||
1382  // vec<RET, NumElements> operatorOP(const vec<DataT, NumElements> &Rhs) const;
1383  // vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
1384 
1385  // OP is: ==, !=, <, >, <=, >=
1386  // vec<RET, NumElements> operatorOP(const vec<DataT, NumElements> &Rhs) const;
1387  // vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
1388 private:
1389  // Generic method that execute "Operation" on underlying values.
1390 
1391 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1392  template <template <typename> class Operation,
1393  typename Ty = vec<DataT, NumElements>>
1394  vec<DataT, NumElements>
1395  operatorHelper(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) const {
1396  vec<DataT, NumElements> Result;
1397  Operation<DataType> Op;
1398  Result.m_Data = Op(m_Data, Rhs.m_Data);
1399  return Result;
1400  }
1401 
1402  template <template <typename> class Operation,
1403  typename Ty = vec<DataT, NumElements>>
1404  vec<DataT, NumElements>
1405  operatorHelper(const EnableIfUsingArrayOnDevice<Ty> &Rhs) const {
1406  vec<DataT, NumElements> Result;
1407  Operation<DataT> Op;
1408  for (size_t I = 0; I < NumElements; ++I) {
1409  Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1410  }
1411  return Result;
1412  }
1413 #else // __SYCL_USE_EXT_VECTOR_TYPE__
1414  template <template <typename> class Operation>
1415  vec<DataT, NumElements>
1416  operatorHelper(const vec<DataT, NumElements> &Rhs) const {
1417  vec<DataT, NumElements> Result;
1418  Operation<DataT> Op;
1419  for (size_t I = 0; I < NumElements; ++I) {
1420  Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1421  }
1422  return Result;
1423  }
1424 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
1425 
1426  // setValue and getValue should be able to operate on different underlying
1427  // types: enum cl_float#N , builtin vector float#N, builtin type float.
1428 
1429 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1430  template <int Num = NumElements, typename Ty = int,
1431  typename = typename std::enable_if_t<1 != Num>>
1432  constexpr void setValue(EnableIfNotHostHalf<Ty> Index, const DataT &Value,
1433  int) {
1434  m_Data[Index] = vec_data<DataT>::get(Value);
1435  }
1436 
1437  template <int Num = NumElements, typename Ty = int,
1438  typename = typename std::enable_if_t<1 != Num>>
1439  constexpr DataT getValue(EnableIfNotHostHalf<Ty> Index, int) const {
1440  return vec_data<DataT>::get(m_Data[Index]);
1441  }
1442 
1443  template <int Num = NumElements, typename Ty = int,
1444  typename = typename std::enable_if_t<1 != Num>>
1445  constexpr void setValue(EnableIfHostHalf<Ty> Index, const DataT &Value, int) {
1446  m_Data.s[Index] = vec_data<DataT>::get(Value);
1447  }
1448 
1449  template <int Num = NumElements, typename Ty = int,
1450  typename = typename std::enable_if_t<1 != Num>>
1451  constexpr DataT getValue(EnableIfHostHalf<Ty> Index, int) const {
1452  return vec_data<DataT>::get(m_Data.s[Index]);
1453  }
1454 #else // __SYCL_USE_EXT_VECTOR_TYPE__
1455  template <int Num = NumElements,
1456  typename = typename std::enable_if_t<1 != Num>>
1457  constexpr void setValue(int Index, const DataT &Value, int) {
1458 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1459  m_Data[Index] = vec_data<DataT>::get(Value);
1460 #else
1461  m_Data.s[Index] = vec_data<DataT>::get(Value);
1462 #endif
1463  }
1464 
1465  template <int Num = NumElements,
1466  typename = typename std::enable_if_t<1 != Num>>
1467  constexpr DataT getValue(int Index, int) const {
1468 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1469  return vec_data<DataT>::get(m_Data[Index]);
1470 #else
1471  return vec_data<DataT>::get(m_Data.s[Index]);
1472 #endif
1473  }
1474 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
1475 
1476  template <int Num = NumElements,
1477  typename = typename std::enable_if_t<1 == Num>>
1478  constexpr void setValue(int, const DataT &Value, float) {
1479  m_Data = vec_data<DataT>::get(Value);
1480  }
1481 
1482  template <int Num = NumElements,
1483  typename = typename std::enable_if_t<1 == Num>>
1484  DataT getValue(int, float) const {
1485  return vec_data<DataT>::get(m_Data);
1486  }
1487 
1488  // Special proxies as specialization is not allowed in class scope.
1489  constexpr void setValue(int Index, const DataT &Value) {
1490  if (NumElements == 1)
1491  setValue(Index, Value, 0);
1492  else
1493  setValue(Index, Value, 0.f);
1494  }
1495 
1496  DataT getValue(int Index) const {
1497  return (NumElements == 1) ? getValue(Index, 0) : getValue(Index, 0.f);
1498  }
1499 
1500 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1501 
1502  // fields
1503 
1504  // Alignment is the same as size, to a maximum size of 64.
1505  // detail::vector_alignment will return that value.
1506  alignas(detail::vector_alignment<DataT, NumElements>::value) DataType m_Data;
1507 
1508 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1509 
1510 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1511 
1512  // fields
1513 
1514  // Used "__SYCL_ALIGNED_VAR" instead "alignas" to handle MSVC compiler.
1515  // For MSVC compiler max alignment is 64, e.g. vec<double, 16> required
1516  // alignment of 128 and MSVC compiler cann't align a parameter with requested
1517  // alignment of 128. For alignment request larger than 64, 64-alignment
1518  // is applied
1519  __SYCL_ALIGNED_VAR(DataType,
1520  (detail::vector_alignment<DataT, NumElements>::value),
1521  m_Data);
1522 
1523 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1524 
1525  // friends
1526  template <typename T1, typename T2, typename T3, template <typename> class T4,
1527  int... T5>
1528  friend class detail::SwizzleOp;
1529  template <typename T1, int T2> friend class vec;
1530 };
1531 
1532 #ifdef __cpp_deduction_guides
1533 // all compilers supporting deduction guides also support fold expressions
1534 template <class T, class... U,
1535  class = std::enable_if_t<(std::is_same_v<T, U> && ...)>>
1536 vec(T, U...) -> vec<T, sizeof...(U) + 1>;
1537 #endif
1538 
1539 namespace detail {
1540 
1541 // SwizzleOP represents expression templates that operate on vec.
1542 // Actual computation performed on conversion or assignment operators.
1543 template <typename VecT, typename OperationLeftT, typename OperationRightT,
1544  template <typename> class OperationCurrentT, int... Indexes>
1545 class SwizzleOp {
1546  using DataT = typename VecT::element_type;
1547  // Certain operators return a vector with a different element type. Also, the
1548  // left and right operand types may differ. CommonDataT selects a result type
1549  // based on these types to ensure that the result value can be represented.
1550  //
1551  // Example 1:
1552  // sycl::vec<unsigned char, 4> vec{...};
1553  // auto result = 300u + vec.x();
1554  //
1555  // CommonDataT is std::common_type_t<OperationLeftT, OperationRightT> since
1556  // it's larger than unsigned char.
1557  //
1558  // Example 2:
1559  // sycl::vec<bool, 1> vec{...};
1560  // auto result = vec.template swizzle<sycl::elem::s0>() && vec;
1561  //
1562  // CommonDataT is DataT since operator&& returns a vector with element type
1563  // int8_t, which is larger than bool.
1564  //
1565  // Example 3:
1566  // sycl::vec<std::byte, 4> vec{...}; auto swlo = vec.lo();
1567  // auto result = swlo == swlo;
1568  //
1569  // CommonDataT is DataT since operator== returns a vector with element type
1570  // int8_t, which is the same size as std::byte. std::common_type_t<DataT, ...>
1571  // can't be used here since there's no type that int8_t and std::byte can both
1572  // be implicitly converted to.
1573  using OpLeftDataT = typename OperationLeftT::DataT;
1574  using OpRightDataT = typename OperationRightT::DataT;
1575  using CommonDataT = std::conditional_t<
1576  sizeof(DataT) >= sizeof(std::common_type_t<OpLeftDataT, OpRightDataT>),
1577  DataT, std::common_type_t<OpLeftDataT, OpRightDataT>>;
1578  static constexpr int getNumElements() { return sizeof...(Indexes); }
1579 
1580  using rel_t = detail::rel_t<DataT>;
1581  using vec_t = vec<DataT, sizeof...(Indexes)>;
1582  using vec_rel_t = vec<rel_t, sizeof...(Indexes)>;
1583 
1584  template <typename OperationRightT_,
1585  template <typename> class OperationCurrentT_, int... Idx_>
1586  using NewLHOp = SwizzleOp<VecT,
1587  SwizzleOp<VecT, OperationLeftT, OperationRightT,
1588  OperationCurrentT, Indexes...>,
1589  OperationRightT_, OperationCurrentT_, Idx_...>;
1590 
1591  template <typename OperationRightT_,
1592  template <typename> class OperationCurrentT_, int... Idx_>
1593  using NewRelOp = SwizzleOp<vec<rel_t, VecT::getNumElements()>,
1594  SwizzleOp<VecT, OperationLeftT, OperationRightT,
1595  OperationCurrentT, Indexes...>,
1596  OperationRightT_, OperationCurrentT_, Idx_...>;
1597 
1598  template <typename OperationLeftT_,
1599  template <typename> class OperationCurrentT_, int... Idx_>
1600  using NewRHOp = SwizzleOp<VecT, OperationLeftT_,
1601  SwizzleOp<VecT, OperationLeftT, OperationRightT,
1602  OperationCurrentT, Indexes...>,
1603  OperationCurrentT_, Idx_...>;
1604 
1605  template <int IdxNum, typename T = void>
1606  using EnableIfOneIndex = typename std::enable_if_t<
1607  1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1608 
1609  template <int IdxNum, typename T = void>
1610  using EnableIfMultipleIndexes = typename std::enable_if_t<
1611  1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1612 
1613  template <typename T>
1614  using EnableIfScalarType = typename std::enable_if_t<
1615  std::is_convertible_v<DataT, T> &&
1616  (std::is_fundamental_v<vec_data_t<T>> ||
1617  std::is_same_v<typename std::remove_const_t<T>, half>)>;
1618 
1619  template <typename T>
1620  using EnableIfNoScalarType = typename std::enable_if_t<
1621  !std::is_convertible_v<DataT, T> ||
1622  !(std::is_fundamental_v<vec_data_t<T>> ||
1623  std::is_same_v<typename std::remove_const_t<T>, half>)>;
1624 
1625  template <int... Indices>
1626  using Swizzle =
1627  SwizzleOp<VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1628 
1629  template <int... Indices>
1630  using ConstSwizzle =
1631  SwizzleOp<const VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1632 
1633 public:
1634  using element_type = DataT;
1635 
1636  const DataT &operator[](int i) const {
1637  std::array<int, getNumElements()> Idxs{Indexes...};
1638  return (*m_Vector)[Idxs[i]];
1639  }
1640 
1641  template <typename _T = VecT>
1642  std::enable_if_t<!std::is_const_v<_T>, DataT> &operator[](int i) {
1643  std::array<int, getNumElements()> Idxs{Indexes...};
1644  return (*m_Vector)[Idxs[i]];
1645  }
1646 
1647  __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1648  size_t get_count() const { return size(); }
1649  size_t size() const noexcept { return getNumElements(); }
1650 
1651  template <int Num = getNumElements()>
1652  __SYCL2020_DEPRECATED(
1653  "get_size() is deprecated, please use byte_size() instead")
1654  size_t get_size() const {
1655  return byte_size<Num>();
1656  }
1657 
1658  template <int Num = getNumElements()> size_t byte_size() const noexcept {
1659  return sizeof(DataT) * (Num == 3 ? 4 : Num);
1660  }
1661 
1662  template <typename T, int IdxNum = getNumElements(),
1663  typename = EnableIfOneIndex<IdxNum>,
1664  typename = EnableIfScalarType<T>>
1665  operator T() const {
1666  return getValue(0);
1667  }
1668 
1669  template <typename T, typename = EnableIfScalarType<T>>
1670  friend NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>
1671  operator*(const T &Lhs, const SwizzleOp &Rhs) {
1672  return NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
1673  Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1674  }
1675 
1676  template <typename T, typename = EnableIfScalarType<T>>
1677  friend NewRHOp<GetScalarOp<T>, std::plus, Indexes...>
1678  operator+(const T &Lhs, const SwizzleOp &Rhs) {
1679  return NewRHOp<GetScalarOp<T>, std::plus, Indexes...>(
1680  Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1681  }
1682 
1683  template <typename T, typename = EnableIfScalarType<T>>
1684  friend NewRHOp<GetScalarOp<T>, std::divides, Indexes...>
1685  operator/(const T &Lhs, const SwizzleOp &Rhs) {
1686  return NewRHOp<GetScalarOp<T>, std::divides, Indexes...>(
1687  Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1688  }
1689 
1690  // TODO: Check that Rhs arg is suitable.
1691 #ifdef __SYCL_OPASSIGN
1692 #error "Undefine __SYCL_OPASSIGN macro."
1693 #endif
1694 #define __SYCL_OPASSIGN(OPASSIGN, OP) \
1695  SwizzleOp &operator OPASSIGN(const DataT & Rhs) { \
1696  operatorHelper<OP>(vec_t(Rhs)); \
1697  return *this; \
1698  } \
1699  template <typename RhsOperation> \
1700  SwizzleOp &operator OPASSIGN(const RhsOperation & Rhs) { \
1701  operatorHelper<OP>(Rhs); \
1702  return *this; \
1703  }
1704 
1705  __SYCL_OPASSIGN(+=, std::plus)
1706  __SYCL_OPASSIGN(-=, std::minus)
1707  __SYCL_OPASSIGN(*=, std::multiplies)
1708  __SYCL_OPASSIGN(/=, std::divides)
1709  __SYCL_OPASSIGN(%=, std::modulus)
1710  __SYCL_OPASSIGN(&=, std::bit_and)
1711  __SYCL_OPASSIGN(|=, std::bit_or)
1712  __SYCL_OPASSIGN(^=, std::bit_xor)
1713  __SYCL_OPASSIGN(>>=, RShift)
1714  __SYCL_OPASSIGN(<<=, LShift)
1715 #undef __SYCL_OPASSIGN
1716 
1717 #ifdef __SYCL_UOP
1718 #error "Undefine __SYCL_UOP macro"
1719 #endif
1720 #define __SYCL_UOP(UOP, OPASSIGN) \
1721  SwizzleOp &operator UOP() { \
1722  *this OPASSIGN static_cast<DataT>(1); \
1723  return *this; \
1724  } \
1725  vec_t operator UOP(int) { \
1726  vec_t Ret = *this; \
1727  *this OPASSIGN static_cast<DataT>(1); \
1728  return Ret; \
1729  }
1730 
1731  __SYCL_UOP(++, +=)
1732  __SYCL_UOP(--, -=)
1733 #undef __SYCL_UOP
1734 
1735  template <typename T = DataT>
1736  typename std::enable_if_t<std::is_integral_v<vec_data_t<T>>, vec_t>
1737  operator~() {
1738  vec_t Tmp = *this;
1739  return ~Tmp;
1740  }
1741 
1742  vec_rel_t operator!() {
1743  vec_t Tmp = *this;
1744  return !Tmp;
1745  }
1746 
1747  vec_t operator+() {
1748  vec_t Tmp = *this;
1749  return +Tmp;
1750  }
1751 
1752  vec_t operator-() {
1753  vec_t Tmp = *this;
1754  return -Tmp;
1755  }
1756 
1757  template <int IdxNum = getNumElements(),
1758  typename = EnableIfMultipleIndexes<IdxNum>>
1759  SwizzleOp &operator=(const vec<DataT, IdxNum> &Rhs) {
1760  std::array<int, IdxNum> Idxs{Indexes...};
1761  for (size_t I = 0; I < Idxs.size(); ++I) {
1762  m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1763  }
1764  return *this;
1765  }
1766 
1767  template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
1768  SwizzleOp &operator=(const DataT &Rhs) {
1769  std::array<int, IdxNum> Idxs{Indexes...};
1770  m_Vector->setValue(Idxs[0], Rhs);
1771  return *this;
1772  }
1773 
1774  template <int IdxNum = getNumElements(),
1775  EnableIfMultipleIndexes<IdxNum, bool> = true>
1776  SwizzleOp &operator=(const DataT &Rhs) {
1777  std::array<int, IdxNum> Idxs{Indexes...};
1778  for (auto Idx : Idxs) {
1779  m_Vector->setValue(Idx, Rhs);
1780  }
1781  return *this;
1782  }
1783 
1784  template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
1785  SwizzleOp &operator=(DataT &&Rhs) {
1786  std::array<int, IdxNum> Idxs{Indexes...};
1787  m_Vector->setValue(Idxs[0], Rhs);
1788  return *this;
1789  }
1790 
1791  template <typename T, typename = EnableIfScalarType<T>>
1792  NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>
1793  operator*(const T &Rhs) const {
1794  return NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
1795  m_Vector, *this, GetScalarOp<T>(Rhs));
1796  }
1797 
1798  template <typename RhsOperation,
1799  typename = EnableIfNoScalarType<RhsOperation>>
1800  NewLHOp<RhsOperation, std::multiplies, Indexes...>
1801  operator*(const RhsOperation &Rhs) const {
1802  return NewLHOp<RhsOperation, std::multiplies, Indexes...>(m_Vector, *this,
1803  Rhs);
1804  }
1805 
1806  template <typename T, typename = EnableIfScalarType<T>>
1807  NewLHOp<GetScalarOp<T>, std::plus, Indexes...> operator+(const T &Rhs) const {
1808  return NewLHOp<GetScalarOp<T>, std::plus, Indexes...>(m_Vector, *this,
1809  GetScalarOp<T>(Rhs));
1810  }
1811 
1812  template <typename RhsOperation,
1813  typename = EnableIfNoScalarType<RhsOperation>>
1814  NewLHOp<RhsOperation, std::plus, Indexes...>
1815  operator+(const RhsOperation &Rhs) const {
1816  return NewLHOp<RhsOperation, std::plus, Indexes...>(m_Vector, *this, Rhs);
1817  }
1818 
1819  template <typename T, typename = EnableIfScalarType<T>>
1820  NewLHOp<GetScalarOp<T>, std::minus, Indexes...>
1821  operator-(const T &Rhs) const {
1822  return NewLHOp<GetScalarOp<T>, std::minus, Indexes...>(m_Vector, *this,
1823  GetScalarOp<T>(Rhs));
1824  }
1825 
1826  template <typename RhsOperation,
1827  typename = EnableIfNoScalarType<RhsOperation>>
1828  NewLHOp<RhsOperation, std::minus, Indexes...>
1829  operator-(const RhsOperation &Rhs) const {
1830  return NewLHOp<RhsOperation, std::minus, Indexes...>(m_Vector, *this, Rhs);
1831  }
1832 
1833  template <typename T, typename = EnableIfScalarType<T>>
1834  NewLHOp<GetScalarOp<T>, std::divides, Indexes...>
1835  operator/(const T &Rhs) const {
1836  return NewLHOp<GetScalarOp<T>, std::divides, Indexes...>(
1837  m_Vector, *this, GetScalarOp<T>(Rhs));
1838  }
1839 
1840  template <typename RhsOperation,
1841  typename = EnableIfNoScalarType<RhsOperation>>
1842  NewLHOp<RhsOperation, std::divides, Indexes...>
1843  operator/(const RhsOperation &Rhs) const {
1844  return NewLHOp<RhsOperation, std::divides, Indexes...>(m_Vector, *this,
1845  Rhs);
1846  }
1847 
1848  template <typename T, typename = EnableIfScalarType<T>>
1849  NewLHOp<GetScalarOp<T>, std::modulus, Indexes...>
1850  operator%(const T &Rhs) const {
1851  return NewLHOp<GetScalarOp<T>, std::modulus, Indexes...>(
1852  m_Vector, *this, GetScalarOp<T>(Rhs));
1853  }
1854 
1855  template <typename RhsOperation,
1856  typename = EnableIfNoScalarType<RhsOperation>>
1857  NewLHOp<RhsOperation, std::modulus, Indexes...>
1858  operator%(const RhsOperation &Rhs) const {
1859  return NewLHOp<RhsOperation, std::modulus, Indexes...>(m_Vector, *this,
1860  Rhs);
1861  }
1862 
1863  template <typename T, typename = EnableIfScalarType<T>>
1864  NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>
1865  operator&(const T &Rhs) const {
1866  return NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>(
1867  m_Vector, *this, GetScalarOp<T>(Rhs));
1868  }
1869 
1870  template <typename RhsOperation,
1871  typename = EnableIfNoScalarType<RhsOperation>>
1872  NewLHOp<RhsOperation, std::bit_and, Indexes...>
1873  operator&(const RhsOperation &Rhs) const {
1874  return NewLHOp<RhsOperation, std::bit_and, Indexes...>(m_Vector, *this,
1875  Rhs);
1876  }
1877 
1878  template <typename T, typename = EnableIfScalarType<T>>
1879  NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>
1880  operator|(const T &Rhs) const {
1881  return NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>(
1882  m_Vector, *this, GetScalarOp<T>(Rhs));
1883  }
1884 
1885  template <typename RhsOperation,
1886  typename = EnableIfNoScalarType<RhsOperation>>
1887  NewLHOp<RhsOperation, std::bit_or, Indexes...>
1888  operator|(const RhsOperation &Rhs) const {
1889  return NewLHOp<RhsOperation, std::bit_or, Indexes...>(m_Vector, *this, Rhs);
1890  }
1891 
1892  template <typename T, typename = EnableIfScalarType<T>>
1893  NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>
1894  operator^(const T &Rhs) const {
1895  return NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>(
1896  m_Vector, *this, GetScalarOp<T>(Rhs));
1897  }
1898 
1899  template <typename RhsOperation,
1900  typename = EnableIfNoScalarType<RhsOperation>>
1901  NewLHOp<RhsOperation, std::bit_xor, Indexes...>
1902  operator^(const RhsOperation &Rhs) const {
1903  return NewLHOp<RhsOperation, std::bit_xor, Indexes...>(m_Vector, *this,
1904  Rhs);
1905  }
1906 
1907  template <typename T, typename = EnableIfScalarType<T>>
1908  NewLHOp<GetScalarOp<T>, RShift, Indexes...> operator>>(const T &Rhs) const {
1909  return NewLHOp<GetScalarOp<T>, RShift, Indexes...>(m_Vector, *this,
1910  GetScalarOp<T>(Rhs));
1911  }
1912 
1913  template <typename RhsOperation,
1914  typename = EnableIfNoScalarType<RhsOperation>>
1915  NewLHOp<RhsOperation, RShift, Indexes...>
1916  operator>>(const RhsOperation &Rhs) const {
1917  return NewLHOp<RhsOperation, RShift, Indexes...>(m_Vector, *this, Rhs);
1918  }
1919 
1920  template <typename T, typename = EnableIfScalarType<T>>
1921  NewLHOp<GetScalarOp<T>, LShift, Indexes...> operator<<(const T &Rhs) const {
1922  return NewLHOp<GetScalarOp<T>, LShift, Indexes...>(m_Vector, *this,
1923  GetScalarOp<T>(Rhs));
1924  }
1925 
1926  template <typename RhsOperation,
1927  typename = EnableIfNoScalarType<RhsOperation>>
1928  NewLHOp<RhsOperation, LShift, Indexes...>
1929  operator<<(const RhsOperation &Rhs) const {
1930  return NewLHOp<RhsOperation, LShift, Indexes...>(m_Vector, *this, Rhs);
1931  }
1932 
1933  template <
1934  typename T1, typename T2, typename T3, template <typename> class T4,
1935  int... T5,
1936  typename = typename std::enable_if_t<sizeof...(T5) == getNumElements()>>
1937  SwizzleOp &operator=(const SwizzleOp<T1, T2, T3, T4, T5...> &Rhs) {
1938  std::array<int, getNumElements()> Idxs{Indexes...};
1939  for (size_t I = 0; I < Idxs.size(); ++I) {
1940  m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1941  }
1942  return *this;
1943  }
1944 
1945  template <
1946  typename T1, typename T2, typename T3, template <typename> class T4,
1947  int... T5,
1948  typename = typename std::enable_if_t<sizeof...(T5) == getNumElements()>>
1949  SwizzleOp &operator=(SwizzleOp<T1, T2, T3, T4, T5...> &&Rhs) {
1950  std::array<int, getNumElements()> Idxs{Indexes...};
1951  for (size_t I = 0; I < Idxs.size(); ++I) {
1952  m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1953  }
1954  return *this;
1955  }
1956 
1957  template <typename T, typename = EnableIfScalarType<T>>
1958  NewRelOp<GetScalarOp<T>, EqualTo, Indexes...> operator==(const T &Rhs) const {
1959  return NewRelOp<GetScalarOp<T>, EqualTo, Indexes...>(NULL, *this,
1960  GetScalarOp<T>(Rhs));
1961  }
1962 
1963  template <typename RhsOperation,
1964  typename = EnableIfNoScalarType<RhsOperation>>
1965  NewRelOp<RhsOperation, EqualTo, Indexes...>
1966  operator==(const RhsOperation &Rhs) const {
1967  return NewRelOp<RhsOperation, EqualTo, Indexes...>(NULL, *this, Rhs);
1968  }
1969 
1970  template <typename T, typename = EnableIfScalarType<T>>
1971  NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>
1972  operator!=(const T &Rhs) const {
1973  return NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>(
1974  NULL, *this, GetScalarOp<T>(Rhs));
1975  }
1976 
1977  template <typename RhsOperation,
1978  typename = EnableIfNoScalarType<RhsOperation>>
1979  NewRelOp<RhsOperation, NotEqualTo, Indexes...>
1980  operator!=(const RhsOperation &Rhs) const {
1981  return NewRelOp<RhsOperation, NotEqualTo, Indexes...>(NULL, *this, Rhs);
1982  }
1983 
1984  template <typename T, typename = EnableIfScalarType<T>>
1985  NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>
1986  operator>=(const T &Rhs) const {
1987  return NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>(
1988  NULL, *this, GetScalarOp<T>(Rhs));
1989  }
1990 
1991  template <typename RhsOperation,
1992  typename = EnableIfNoScalarType<RhsOperation>>
1993  NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>
1994  operator>=(const RhsOperation &Rhs) const {
1995  return NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>(NULL, *this, Rhs);
1996  }
1997 
1998  template <typename T, typename = EnableIfScalarType<T>>
1999  NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>
2000  operator<=(const T &Rhs) const {
2001  return NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>(
2002  NULL, *this, GetScalarOp<T>(Rhs));
2003  }
2004 
2005  template <typename RhsOperation,
2006  typename = EnableIfNoScalarType<RhsOperation>>
2007  NewRelOp<RhsOperation, LessEqualTo, Indexes...>
2008  operator<=(const RhsOperation &Rhs) const {
2009  return NewRelOp<RhsOperation, LessEqualTo, Indexes...>(NULL, *this, Rhs);
2010  }
2011 
2012  template <typename T, typename = EnableIfScalarType<T>>
2013  NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>
2014  operator>(const T &Rhs) const {
2015  return NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>(
2016  NULL, *this, GetScalarOp<T>(Rhs));
2017  }
2018 
2019  template <typename RhsOperation,
2020  typename = EnableIfNoScalarType<RhsOperation>>
2021  NewRelOp<RhsOperation, GreaterThan, Indexes...>
2022  operator>(const RhsOperation &Rhs) const {
2023  return NewRelOp<RhsOperation, GreaterThan, Indexes...>(NULL, *this, Rhs);
2024  }
2025 
2026  template <typename T, typename = EnableIfScalarType<T>>
2027  NewRelOp<GetScalarOp<T>, LessThan, Indexes...> operator<(const T &Rhs) const {
2028  return NewRelOp<GetScalarOp<T>, LessThan, Indexes...>(NULL, *this,
2029  GetScalarOp<T>(Rhs));
2030  }
2031 
2032  template <typename RhsOperation,
2033  typename = EnableIfNoScalarType<RhsOperation>>
2034  NewRelOp<RhsOperation, LessThan, Indexes...>
2035  operator<(const RhsOperation &Rhs) const {
2036  return NewRelOp<RhsOperation, LessThan, Indexes...>(NULL, *this, Rhs);
2037  }
2038 
2039  template <typename T, typename = EnableIfScalarType<T>>
2040  NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>
2041  operator&&(const T &Rhs) const {
2042  return NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>(
2043  NULL, *this, GetScalarOp<T>(Rhs));
2044  }
2045 
2046  template <typename RhsOperation,
2047  typename = EnableIfNoScalarType<RhsOperation>>
2048  NewRelOp<RhsOperation, LogicalAnd, Indexes...>
2049  operator&&(const RhsOperation &Rhs) const {
2050  return NewRelOp<RhsOperation, LogicalAnd, Indexes...>(NULL, *this, Rhs);
2051  }
2052 
2053  template <typename T, typename = EnableIfScalarType<T>>
2054  NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>
2055  operator||(const T &Rhs) const {
2056  return NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>(NULL, *this,
2057  GetScalarOp<T>(Rhs));
2058  }
2059 
2060  template <typename RhsOperation,
2061  typename = EnableIfNoScalarType<RhsOperation>>
2062  NewRelOp<RhsOperation, LogicalOr, Indexes...>
2063  operator||(const RhsOperation &Rhs) const {
2064  return NewRelOp<RhsOperation, LogicalOr, Indexes...>(NULL, *this, Rhs);
2065  }
2066 
2067  // Begin hi/lo, even/odd, xyzw, and rgba swizzles.
2068 private:
2069  // Indexer used in the swizzles.def.
2070  // Currently it is defined as a template struct. Replacing it with a constexpr
2071  // function would activate a bug in MSVC that is fixed only in v19.20.
2072  // Until then MSVC does not recognize such constexpr functions as const and
2073  // thus does not let using them in template parameters inside swizzle.def.
2074  template <int Index> struct Indexer {
2075  static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...};
2076  static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index];
2077  };
2078 
2079 public:
2080 #ifdef __SYCL_ACCESS_RETURN
2081 #error "Undefine __SYCL_ACCESS_RETURN macro"
2082 #endif
2083 #define __SYCL_ACCESS_RETURN m_Vector
2084 #include "swizzles.def"
2085 #undef __SYCL_ACCESS_RETURN
2086  // End of hi/lo, even/odd, xyzw, and rgba swizzles.
2087 
2088  // Leave store() interface to automatic conversion to vec<>.
2089  // Load to vec_t and then assign to swizzle.
2090  template <access::address_space Space, access::decorated DecorateAddress>
2091  void load(size_t offset, multi_ptr<DataT, Space, DecorateAddress> ptr) {
2092  vec_t Tmp;
2093  Tmp.template load(offset, ptr);
2094  *this = Tmp;
2095  }
2096 
2097  template <typename convertT, rounding_mode roundingMode>
2098  vec<convertT, sizeof...(Indexes)> convert() const {
2099  // First materialize the swizzle to vec_t and then apply convert() to it.
2100  vec_t Tmp = *this;
2101  return Tmp.template convert<convertT, roundingMode>();
2102  }
2103 
2104  template <typename asT> asT as() const {
2105  // First materialize the swizzle to vec_t and then apply as() to it.
2106  vec_t Tmp = *this;
2107  static_assert((sizeof(Tmp) == sizeof(asT)),
2108  "The new SYCL vec type must have the same storage size in "
2109  "bytes as this SYCL swizzled vec");
2110  static_assert(
2111  detail::is_contained<asT, detail::gtl::vector_basic_list>::value ||
2112  detail::is_contained<asT, detail::gtl::vector_bool_list>::value,
2113  "asT must be SYCL vec of a different element type and "
2114  "number of elements specified by asT");
2115  return Tmp.template as<asT>();
2116  }
2117 
2118 private:
2119  SwizzleOp(const SwizzleOp &Rhs)
2120  : m_Vector(Rhs.m_Vector), m_LeftOperation(Rhs.m_LeftOperation),
2121  m_RightOperation(Rhs.m_RightOperation) {}
2122 
2123  SwizzleOp(VecT *Vector, OperationLeftT LeftOperation,
2124  OperationRightT RightOperation)
2125  : m_Vector(Vector), m_LeftOperation(LeftOperation),
2126  m_RightOperation(RightOperation) {}
2127 
2128  SwizzleOp(VecT *Vector) : m_Vector(Vector) {}
2129 
2130  SwizzleOp(SwizzleOp &&Rhs)
2131  : m_Vector(Rhs.m_Vector), m_LeftOperation(std::move(Rhs.m_LeftOperation)),
2132  m_RightOperation(std::move(Rhs.m_RightOperation)) {}
2133 
2134  // Either performing CurrentOperation on results of left and right operands
2135  // or reading values from actual vector. Perform implicit type conversion when
2136  // the number of elements == 1
2137 
2138  template <int IdxNum = getNumElements()>
2139  CommonDataT getValue(EnableIfOneIndex<IdxNum, size_t> Index) const {
2140  if (std::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
2141  std::array<int, getNumElements()> Idxs{Indexes...};
2142  return m_Vector->getValue(Idxs[Index]);
2143  }
2144  auto Op = OperationCurrentT<vec_data_t<CommonDataT>>();
2145  return vec_data<CommonDataT>::get(
2146  Op(vec_data<CommonDataT>::get(m_LeftOperation.getValue(Index)),
2147  vec_data<CommonDataT>::get(m_RightOperation.getValue(Index))));
2148  }
2149 
2150  template <int IdxNum = getNumElements()>
2151  DataT getValue(EnableIfMultipleIndexes<IdxNum, size_t> Index) const {
2152  if (std::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
2153  std::array<int, getNumElements()> Idxs{Indexes...};
2154  return m_Vector->getValue(Idxs[Index]);
2155  }
2156  auto Op = OperationCurrentT<vec_data_t<DataT>>();
2157  return vec_data<DataT>::get(
2158  Op(vec_data<DataT>::get(m_LeftOperation.getValue(Index)),
2159  vec_data<DataT>::get(m_RightOperation.getValue(Index))));
2160  }
2161 
2162  template <template <typename> class Operation, typename RhsOperation>
2163  void operatorHelper(const RhsOperation &Rhs) {
2164  Operation<vec_data_t<DataT>> Op;
2165  std::array<int, getNumElements()> Idxs{Indexes...};
2166  for (size_t I = 0; I < Idxs.size(); ++I) {
2167  DataT Res = vec_data<DataT>::get(
2168  Op(vec_data<DataT>::get(m_Vector->getValue(Idxs[I])),
2169  vec_data<DataT>::get(Rhs.getValue(I))));
2170  m_Vector->setValue(Idxs[I], Res);
2171  }
2172  }
2173 
2174  // fields
2175  VecT *m_Vector;
2176 
2177  OperationLeftT m_LeftOperation;
2178  OperationRightT m_RightOperation;
2179 
2180  // friends
2181  template <typename T1, int T2> friend class sycl::vec;
2182 
2183  template <typename T1, typename T2, typename T3, template <typename> class T4,
2184  int... T5>
2185  friend class SwizzleOp;
2186 };
2187 } // namespace detail
2188 
2189 // scalar BINOP vec<>
2190 // scalar BINOP SwizzleOp
2191 // vec<> BINOP SwizzleOp
2192 #ifdef __SYCL_BINOP
2193 #error "Undefine __SYCL_BINOP macro"
2194 #endif
2195 #define __SYCL_BINOP(BINOP) \
2196  template <typename T, int Num> \
2197  typename std::enable_if_t< \
2198  std::is_fundamental_v<vec_data_t<T>> || \
2199  std::is_same_v<typename std::remove_const_t<T>, half>, \
2200  vec<T, Num>> \
2201  operator BINOP(const T & Lhs, const vec<T, Num> &Rhs) { \
2202  return vec<T, Num>(Lhs) BINOP Rhs; \
2203  } \
2204  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2205  template <typename> class OperationCurrentT, int... Indexes, \
2206  typename T, typename T1 = typename VecT::element_type, \
2207  int Num = sizeof...(Indexes)> \
2208  typename std::enable_if_t< \
2209  std::is_convertible_v<T, T1> && \
2210  (std::is_fundamental_v<vec_data_t<T>> || \
2211  std::is_same_v<typename std::remove_const_t<T>, half>), \
2212  vec<T1, Num>> \
2213  operator BINOP( \
2214  const T & Lhs, \
2215  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2216  OperationCurrentT, Indexes...> &Rhs) { \
2217  vec<T1, Num> Tmp = Rhs; \
2218  return Lhs BINOP Tmp; \
2219  } \
2220  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2221  template <typename> class OperationCurrentT, int... Indexes, \
2222  typename T = typename VecT::element_type, \
2223  int Num = sizeof...(Indexes)> \
2224  vec<T, Num> operator BINOP( \
2225  const vec<T, Num> &Lhs, \
2226  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2227  OperationCurrentT, Indexes...> &Rhs) { \
2228  vec<T, Num> Tmp = Rhs; \
2229  return Lhs BINOP Tmp; \
2230  }
2231 
2232 __SYCL_BINOP(+)
2233 __SYCL_BINOP(-)
2234 __SYCL_BINOP(*)
2235 __SYCL_BINOP(/)
2236 __SYCL_BINOP(%)
2237 __SYCL_BINOP(&)
2238 __SYCL_BINOP(|)
2239 __SYCL_BINOP(^)
2240 __SYCL_BINOP(>>)
2241 __SYCL_BINOP(<<)
2242 #undef __SYCL_BINOP
2243 
2244 // scalar RELLOGOP vec<>
2245 // scalar RELLOGOP SwizzleOp
2246 // vec<> RELLOGOP SwizzleOp
2247 #ifdef __SYCL_RELLOGOP
2248 #error "Undefine __SYCL_RELLOGOP macro"
2249 #endif
2250 #define __SYCL_RELLOGOP(RELLOGOP) \
2251  template <typename T, typename DataT, int Num> \
2252  typename std::enable_if_t< \
2253  std::is_convertible_v<T, DataT> && \
2254  (std::is_fundamental_v<vec_data_t<T>> || \
2255  std::is_same_v<typename std::remove_const_t<T>, half>), \
2256  vec<detail::rel_t<DataT>, Num>> \
2257  operator RELLOGOP(const T & Lhs, const vec<DataT, Num> &Rhs) { \
2258  return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs; \
2259  } \
2260  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2261  template <typename> class OperationCurrentT, int... Indexes, \
2262  typename T, typename T1 = typename VecT::element_type, \
2263  int Num = sizeof...(Indexes)> \
2264  typename std::enable_if_t< \
2265  std::is_convertible_v<T, T1> && \
2266  (std::is_fundamental_v<vec_data_t<T>> || \
2267  std::is_same_v<typename std::remove_const_t<T>, half>), \
2268  vec<detail::rel_t<T1>, Num>> \
2269  operator RELLOGOP( \
2270  const T & Lhs, \
2271  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2272  OperationCurrentT, Indexes...> &Rhs) { \
2273  vec<T1, Num> Tmp = Rhs; \
2274  return Lhs RELLOGOP Tmp; \
2275  } \
2276  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2277  template <typename> class OperationCurrentT, int... Indexes, \
2278  typename T = typename VecT::element_type, \
2279  int Num = sizeof...(Indexes)> \
2280  vec<detail::rel_t<T>, Num> operator RELLOGOP( \
2281  const vec<T, Num> &Lhs, \
2282  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2283  OperationCurrentT, Indexes...> &Rhs) { \
2284  vec<T, Num> Tmp = Rhs; \
2285  return Lhs RELLOGOP Tmp; \
2286  }
2287 
2288 __SYCL_RELLOGOP(==)
2289 __SYCL_RELLOGOP(!=)
2290 __SYCL_RELLOGOP(>)
2291 __SYCL_RELLOGOP(<)
2292 __SYCL_RELLOGOP(>=)
2293 __SYCL_RELLOGOP(<=)
2294 // TODO: limit to integral types.
2295 __SYCL_RELLOGOP(&&)
2296 __SYCL_RELLOGOP(||)
2297 #undef __SYCL_RELLOGOP
2298 
2299 namespace detail {
2300 
2301 // Vectors of size 1 are handled separately and therefore 1 is not included in
2302 // the check below.
2303 constexpr bool isValidVectorSize(int N) {
2304  return N == 2 || N == 3 || N == 4 || N == 8 || N == 16;
2305 }
2306 template <typename T, int N, typename V> struct VecStorage {
2307  static_assert(
2308  isValidVectorSize(N) || N == 1,
2309  "Incorrect number of elements for sycl::vec: only 1, 2, 3, 4, 8 "
2310  "or 16 are supported");
2311  static_assert(!std::is_same_v<V, void>, "Incorrect data type for sycl::vec");
2312 };
2313 
2314 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2315 
2316 #ifdef __SYCL_DEVICE_ONLY__
2317 // device always has ext vector support, but for huge vectors
2318 // we switch to std::array, so that we can use a smaller alignment (64)
2319 // this is to support MSVC, which has a max of 64 for direct params.
2320 template <typename T, int N> struct VecStorageImpl {
2321  static constexpr size_t Num = (N == 3) ? 4 : N;
2322  static constexpr size_t Sz = Num * sizeof(T);
2323  using DataType =
2324  typename std::conditional<Sz <= 64, T __attribute__((ext_vector_type(N))),
2325  std::array<T, Num>>::type;
2326  using VectorDataType = T __attribute__((ext_vector_type(N)));
2327 };
2328 #else // __SYCL_DEVICE_ONLY__
2329 template <typename T, int N> struct VecStorageImpl {
2330  using DataType = std::array<T, (N == 3) ? 4 : N>;
2331 };
2332 #endif // __SYCL_DEVICE_ONLY__
2333 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2334 
2335 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2336 
2337 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
2338 template <typename T, int N> struct VecStorageImpl {
2339  using DataType = T __attribute__((ext_vector_type(N)));
2340 };
2341 #else
2342 // When ext_vector_type is not available, we rely on cl_* types from CL/cl.h
2343 // to represent vec storage.
2344 template <typename T, int N> struct VecStorageImpl;
2345 #define __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, num) \
2346  template <> struct VecStorageImpl<type, num> { \
2347  using DataType = ::cl_##cl_type##num; \
2348  };
2349 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
2350 
2351 #ifndef __SYCL_USE_EXT_VECTOR_TYPE__
2352 #define __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(type, cl_type) \
2353  __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 2) \
2354  __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 3) \
2355  __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 4) \
2356  __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 8) \
2357  __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 16)
2358 
2359 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int8_t, char)
2360 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int16_t, short)
2361 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int32_t, int)
2362 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int64_t, long)
2363 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint8_t, uchar)
2364 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint16_t, ushort)
2365 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint32_t, uint)
2366 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint64_t, ulong)
2367 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(float, float)
2368 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(double, double)
2369 
2370 #undef __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE
2371 #undef __SYCL_DEFINE_VECSTORAGE_IMPL
2372 #endif // ifndef __SYCL_USE_EXT_VECTOR_TYPE__
2373 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2374 
2375 // Single element bool
2376 template <> struct VecStorage<bool, 1, void> {
2377  using DataType = bool;
2378 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2379 #ifdef __SYCL_DEVICE_ONLY__
2380  using VectorDataType = bool;
2381 #endif // __SYCL_DEVICE_ONLY__
2382 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2383 };
2384 
2385 // Multiple element bool
2386 template <int N>
2387 struct VecStorage<bool, N, typename std::enable_if_t<isValidVectorSize(N)>> {
2388  using DataType =
2389  typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
2390  std::int32_t, std::int64_t>,
2391  N>::DataType;
2392 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2393 #ifdef __SYCL_DEVICE_ONLY__
2394  using VectorDataType =
2395  typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
2396  std::int32_t, std::int64_t>,
2397  N>::VectorDataType;
2398 #endif // __SYCL_DEVICE_ONLY__
2399 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2400 };
2401 
2402 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2403 // Single element byte. Multiple elements will propagate through a later
2404 // specialization.
2405 template <> struct VecStorage<std::byte, 1, void> {
2406  using DataType = std::int8_t;
2407 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2408 #ifdef __SYCL_DEVICE_ONLY__
2409  using VectorDataType = std::int8_t;
2410 #endif // __SYCL_DEVICE_ONLY__
2411 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2412 };
2413 #endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2414 
2415 // Single element signed integers
2416 template <typename T>
2417 struct VecStorage<T, 1, typename std::enable_if_t<is_sigeninteger_v<T>>> {
2418 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
2419  using DataType = select_apply_cl_t<T, std::int8_t, std::int16_t, std::int32_t,
2420  std::int64_t>;
2421 #else // __INTEL_PREVIEW_BREAKING_CHANGES
2422  using DataType = T;
2423 #ifdef __SYCL_DEVICE_ONLY__
2424  using VectorDataType = DataType;
2425 #endif // __SYCL_DEVICE_ONLY__
2426 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2427 };
2428 
2429 // Single element unsigned integers
2430 template <typename T>
2431 struct VecStorage<T, 1, typename std::enable_if_t<is_sugeninteger_v<T>>> {
2432 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
2433  using DataType = select_apply_cl_t<T, std::uint8_t, std::uint16_t,
2434  std::uint32_t, std::uint64_t>;
2435 #else // __INTEL_PREVIEW_BREAKING_CHANGES
2436  using DataType = T;
2437 #ifdef __SYCL_DEVICE_ONLY__
2438  using VectorDataType = DataType;
2439 #endif // __SYCL_DEVICE_ONLY__
2440 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2441 };
2442 
2443 // Single element floating-point (except half)
2444 template <typename T>
2445 struct VecStorage<
2446  T, 1, typename std::enable_if_t<!is_half_v<T> && is_sgenfloat_v<T>>> {
2447  using DataType = T;
2448 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2449 #ifdef __SYCL_DEVICE_ONLY__
2450  using VectorDataType = DataType;
2451 #endif // __SYCL_DEVICE_ONLY__
2452 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2453 };
2454 // Multiple elements signed/unsigned integers and floating-point (except half)
2455 template <typename T, int N>
2456 struct VecStorage<
2457  T, N,
2458  typename std::enable_if_t<isValidVectorSize(N) &&
2459  (is_sgeninteger_v<T> ||
2460  (is_sgenfloat_v<T> && !is_half_v<T>))>> {
2461  using DataType =
2462  typename VecStorageImpl<typename VecStorage<T, 1>::DataType, N>::DataType;
2463 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2464 #ifdef __SYCL_DEVICE_ONLY__
2465  using VectorDataType =
2466  typename VecStorageImpl<typename VecStorage<T, 1>::DataType,
2467  N>::VectorDataType;
2468 #endif // __SYCL_DEVICE_ONLY__
2469 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2470 };
2471 
2472 // Single element half
2473 template <> struct VecStorage<half, 1, void> {
2474  using DataType = sycl::detail::half_impl::StorageT;
2475 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2476 #ifdef __SYCL_DEVICE_ONLY__
2477  using VectorDataType = sycl::detail::half_impl::StorageT;
2478 #endif // __SYCL_DEVICE_ONLY__
2479 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
2480 };
2481 
2482 // Multiple elements half
2483 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
2484 #define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
2485  template <> struct VecStorage<half, Num, void> { \
2486  using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
2487  using VectorDataType = sycl::detail::half_impl::Vec##Num##StorageT; \
2488  };
2489 #else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) &&
2490  // defined(__SYCL_DEVICE_ONLY__)
2491 #define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
2492  template <> struct VecStorage<half, Num, void> { \
2493  using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
2494  };
2495 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) &&
2496  // defined(__SYCL_DEVICE_ONLY__)
2497 
2498 __SYCL_DEFINE_HALF_VECSTORAGE(2)
2499 __SYCL_DEFINE_HALF_VECSTORAGE(3)
2500 __SYCL_DEFINE_HALF_VECSTORAGE(4)
2501 __SYCL_DEFINE_HALF_VECSTORAGE(8)
2502 __SYCL_DEFINE_HALF_VECSTORAGE(16)
2503 #undef __SYCL_DEFINE_HALF_VECSTORAGE
2504 } // namespace detail
2505 
2509 #define SYCL_DEVICE_COPYABLE 1
2510 
2518 template <typename T> struct is_device_copyable;
2519 
2520 namespace detail {
2521 template <typename T, typename = void>
2522 struct is_device_copyable_impl : std::is_trivially_copyable<T> {};
2523 
2524 template <typename T>
2525 struct is_device_copyable_impl<
2526  T, std::enable_if_t<!std::is_same_v<T, std::remove_cv_t<T>>>>
2527  // Cannot express this "recursion" (to take user's partial non-cv
2528  // specializations into account) without this helper struct.
2529  : is_device_copyable<std::remove_cv_t<T>> {};
2530 } // namespace detail
2531 
2532 template <typename T>
2533 struct is_device_copyable : detail::is_device_copyable_impl<T> {};
2534 
2535 // std::array<T, 0> is implicitly device copyable type.
2536 template <typename T>
2537 struct is_device_copyable<std::array<T, 0>> : std::true_type {};
2538 
2539 // std::array<T, N> is implicitly device copyable type if T is device copyable.
2540 template <typename T, std::size_t N>
2541 struct is_device_copyable<std::array<T, N>> : is_device_copyable<T> {};
2542 
2543 // std::optional<T> is implicitly device copyable type if T is device copyable.
2544 template <typename T>
2545 struct is_device_copyable<std::optional<T>> : is_device_copyable<T> {};
2546 
2547 // std::pair<T1, T2> is implicitly device copyable type if T1 and T2 are device
2548 // copyable.
2549 template <typename T1, typename T2>
2550 struct is_device_copyable<std::pair<T1, T2>>
2551  : std::bool_constant<is_device_copyable<T1>::value &&
2552  is_device_copyable<T2>::value> {};
2553 
2554 // std::tuple<Ts...> is implicitly device copyable type if each type T of Ts...
2555 // is device copyable.
2556 template <typename... Ts>
2557 struct is_device_copyable<std::tuple<Ts...>>
2558  : std::bool_constant<(... && is_device_copyable<Ts>::value)> {};
2559 
2560 // std::variant<Ts...> is implicitly device copyable type if each type T of
2561 // Ts... is device copyable.
2562 template <typename... Ts>
2563 struct is_device_copyable<std::variant<Ts...>>
2564  : std::bool_constant<(... && is_device_copyable<Ts>::value)> {};
2565 
2566 // marray is device copyable if element type is device copyable.
2567 template <typename T, std::size_t N>
2568 struct is_device_copyable<sycl::marray<T, N>> : is_device_copyable<T> {};
2569 
2570 // array is device copyable if element type is device copyable.
2571 template <typename T, std::size_t N>
2572 struct is_device_copyable<T[N]> : is_device_copyable<T> {};
2573 
2574 template <typename T>
2575 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
2576 
2577 namespace detail {
2578 template <typename T, typename = void>
2579 struct IsDeprecatedDeviceCopyable : std::false_type {};
2580 
2581 // TODO: using C++ attribute [[deprecated]] or the macro __SYCL2020_DEPRECATED
2582 // does not produce expected warning message for the type 'T'.
2583 template <typename T>
2584 struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
2585  IsDeprecatedDeviceCopyable<
2586  T, std::enable_if_t<std::is_trivially_copy_constructible_v<T> &&
2587  std::is_trivially_destructible_v<T> &&
2588  !is_device_copyable_v<T>>> : std::true_type {};
2589 
2590 template <typename T, int N>
2591 struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
2592  IsDeprecatedDeviceCopyable<T[N]> : IsDeprecatedDeviceCopyable<T> {};
2593 
2594 #ifdef __SYCL_DEVICE_ONLY__
2595 // Checks that the fields of the type T with indices 0 to (NumFieldsToCheck -
2596 // 1) are device copyable.
2597 template <typename T, unsigned NumFieldsToCheck>
2598 struct CheckFieldsAreDeviceCopyable
2599  : CheckFieldsAreDeviceCopyable<T, NumFieldsToCheck - 1> {
2600  using FieldT = decltype(__builtin_field_type(T, NumFieldsToCheck - 1));
2601  static_assert(is_device_copyable_v<FieldT> ||
2602  detail::IsDeprecatedDeviceCopyable<FieldT>::value,
2603  "The specified type is not device copyable");
2604 };
2605 
2606 template <typename T> struct CheckFieldsAreDeviceCopyable<T, 0> {};
2607 
2608 // Checks that the base classes of the type T with indices 0 to
2609 // (NumFieldsToCheck - 1) are device copyable.
2610 template <typename T, unsigned NumBasesToCheck>
2611 struct CheckBasesAreDeviceCopyable
2612  : CheckBasesAreDeviceCopyable<T, NumBasesToCheck - 1> {
2613  using BaseT = decltype(__builtin_base_type(T, NumBasesToCheck - 1));
2614  static_assert(is_device_copyable_v<BaseT> ||
2615  detail::IsDeprecatedDeviceCopyable<BaseT>::value,
2616  "The specified type is not device copyable");
2617 };
2618 
2619 template <typename T> struct CheckBasesAreDeviceCopyable<T, 0> {};
2620 
2621 // All the captures of a lambda or functor of type FuncT passed to a kernel
2622 // must be is_device_copyable, which extends to bases and fields of FuncT.
2623 // Fields are captures of lambda/functors and bases are possible base classes
2624 // of functors also allowed by SYCL.
2625 // The SYCL-2020 implementation must check each of the fields & bases of the
2626 // type FuncT, only one level deep, which is enough to see if they are all
2627 // device copyable by using the result of is_device_copyable returned for them.
2628 // At this moment though the check also allowes using types for which
2629 // (is_trivially_copy_constructible && is_trivially_destructible) returns true
2630 // and (is_device_copyable) returns false. That is the deprecated behavior and
2631 // is currently/temporarily supported only to not break older SYCL programs.
2632 template <typename FuncT>
2633 struct CheckDeviceCopyable
2634  : CheckFieldsAreDeviceCopyable<FuncT, __builtin_num_fields(FuncT)>,
2635  CheckBasesAreDeviceCopyable<FuncT, __builtin_num_bases(FuncT)> {};
2636 
2637 // Below are two specializations for CheckDeviceCopyable when a kernel lambda
2638 // is wrapped after range rounding optimization.
2639 template <typename TransformedArgType, int Dims, typename KernelType>
2640 struct CheckDeviceCopyable<
2641  RoundedRangeKernel<TransformedArgType, Dims, KernelType>>
2642  : CheckDeviceCopyable<KernelType> {};
2643 
2644 template <typename TransformedArgType, int Dims, typename KernelType>
2645 struct CheckDeviceCopyable<
2646  RoundedRangeKernelWithKH<TransformedArgType, Dims, KernelType>>
2647  : CheckDeviceCopyable<KernelType> {};
2648 
2649 #endif // __SYCL_DEVICE_ONLY__
2650 } // namespace detail
2651 
2652 } // namespace _V1
2653 } // namespace sycl
2654 
2655 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2656 #undef __SYCL_ALIGNED_VAR
2657 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
sycl::_V1::elem::g
static constexpr int g
Definition: types.hpp:112
sycl::_V1::elem::sC
static constexpr int sC
Definition: types.hpp:127
sycl::_V1::detail::GetScalarOp::GetScalarOp
GetScalarOp(DataT Data)
Definition: types.hpp:195
sycl::_V1::opencl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:134
sycl::_V1::detail::LogicalOr
Definition: types.hpp:244
sycl::_V1::elem::b
static constexpr int b
Definition: types.hpp:113
sycl::_V1::detail::GreaterThan
Definition: types.hpp:226
sycl::_V1::elem::s4
static constexpr int s4
Definition: types.hpp:119
type_traits.hpp
sycl::_V1::elem
Definition: types.hpp:106
sycl::_V1::detail::GreaterThan::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:227
memcpy.hpp
sycl::_V1::detail::GreaterEqualTo
Definition: types.hpp:214
sycl::_V1::detail::host_half_impl::half
Definition: half_type.hpp:150
sycl::_V1::elem::s7
static constexpr int s7
Definition: types.hpp:122
sycl::_V1::detail::EqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:203
sycl::_V1::elem::y
static constexpr int y
Definition: types.hpp:108
sycl::_V1::elem::sE
static constexpr int sE
Definition: types.hpp:129
sycl::_V1::detail::int64_t
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:35
sycl::_V1::detail::LogicalAnd
Definition: types.hpp:238
sycl::_V1::detail::vec_helper::get
static constexpr RetType get(T value)
Definition: types.hpp:144
sycl::_V1::elem::s9
static constexpr int s9
Definition: types.hpp:124
sycl::_V1::elem::s8
static constexpr int s8
Definition: types.hpp:123
sycl::_V1::elem::s0
static constexpr int s0
Definition: types.hpp:115
sycl::_V1::detail::vec_helper< std::byte >::get
static constexpr std::byte get(std::uint8_t value)
Definition: types.hpp:156
sycl::_V1::vec_data_t
typename detail::vec_helper< T >::RetType vec_data_t
Definition: types.hpp:273
sycl::_V1::byte
unsigned char byte
Definition: image.hpp:108
sycl::_V1::detail::RShift::operator()
constexpr T operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:251
sycl
Definition: access.hpp:18
sycl::_V1::detail::NotEqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:209
access.hpp
sycl::_V1::opencl::cl_short
std::int16_t cl_short
Definition: aliases.hpp:132
sycl::_V1::detail::select_apply_cl_t
std::conditional_t< sizeof(_IN)==1, T8, std::conditional_t< sizeof(_IN)==2, T16, std::conditional_t< sizeof(_IN)==4, T32, T64 > >> select_apply_cl_t
Definition: types.hpp:140
vector_convert.hpp
generic_type_lists.hpp
sycl::_V1::elem::a
static constexpr int a
Definition: types.hpp:114
sycl::_V1::detail::vec_helper< bool >::get
static constexpr RetType get(bool value)
Definition: types.hpp:149
bool
sycl::_V1::elem::s5
static constexpr int s5
Definition: types.hpp:120
sycl::_V1::detail::vec_helper::RetType
T RetType
Definition: types.hpp:143
sycl::_V1::detail::RShift
Definition: types.hpp:250
defines_elementary.hpp
sycl::_V1::elem::z
static constexpr int z
Definition: types.hpp:109
sycl::_V1::opencl::cl_char
std::int8_t cl_char
Definition: aliases.hpp:130
sycl::_V1::elem::s3
static constexpr int s3
Definition: types.hpp:118
sycl::_V1::detail::vec_helper< bool >::RetType
select_apply_cl_t< bool, std::int8_t, std::int16_t, std::int32_t, std::int64_t > RetType
Definition: types.hpp:148
cl.h
sycl::_V1::elem::sA
static constexpr int sA
Definition: types.hpp:125
generic_type_traits.hpp
sycl::_V1::detail::vec_helper< std::byte >::get
static constexpr RetType get(std::byte value)
Definition: types.hpp:155
sycl::_V1::opencl::cl_long
std::int64_t cl_long
Definition: aliases.hpp:136
multi_ptr.hpp
sycl::_V1::elem::sB
static constexpr int sB
Definition: types.hpp:126
sycl::_V1::elem::s2
static constexpr int s2
Definition: types.hpp:117
common.hpp
sycl::_V1::detail::LShift
Definition: types.hpp:256
sycl::_V1::elem::sF
static constexpr int sF
Definition: types.hpp:130
sycl::_V1::detail::LessEqualTo
Definition: types.hpp:220
sycl::_V1::elem::r
static constexpr int r
Definition: types.hpp:111
sycl::_V1::elem::s1
static constexpr int s1
Definition: types.hpp:116
sycl::_V1::detail::GetScalarOp
Definition: types.hpp:192
sycl::_V1::detail::rel_t
typename std::conditional_t< sizeof(DataT)==sizeof(opencl::cl_char), opencl::cl_char, typename std::conditional_t< sizeof(DataT)==sizeof(opencl::cl_short), opencl::cl_short, typename std::conditional_t< sizeof(DataT)==sizeof(opencl::cl_int), opencl::cl_int, typename std::conditional_t< sizeof(DataT)==sizeof(opencl::cl_long), opencl::cl_long, bool > >> > rel_t
Definition: types.hpp:178
sycl::_V1::detail::EqualTo
Definition: types.hpp:202
sycl::_V1::detail::LShift::operator()
constexpr T operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:257
sycl::_V1::detail::GetScalarOp::getValue
DataT getValue(size_t) const
Definition: types.hpp:196
sycl::_V1::elem::x
static constexpr int x
Definition: types.hpp:107
sycl::_V1::detail::LogicalAnd::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:239
exception.hpp
std
Definition: accessor.hpp:4171
sycl::_V1::detail::VecStorage
Definition: types.hpp:166
sycl::_V1::detail::GetOp::operator()
DataT operator()(DataT, DataT)
Definition: types.hpp:186
half_type.hpp
sycl::_V1::elem::w
static constexpr int w
Definition: types.hpp:110
sycl::_V1::detail::LogicalOr::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:245
aliases.hpp
marray.hpp
sycl::_V1::detail::GreaterEqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:215
sycl::_V1::detail::half_impl::StorageT
detail::host_half_impl::half StorageT
Definition: half_type.hpp:253
sycl::_V1::elem::sD
static constexpr int sD
Definition: types.hpp:128
vector_traits.hpp
sycl::_V1::elem::s6
static constexpr int s6
Definition: types.hpp:121
sycl::_V1::detail::GetOp
Definition: types.hpp:182
sycl::_V1::detail::GetOp::getValue
DataT getValue(size_t) const
Definition: types.hpp:185
sycl::_V1::detail::GetOp::DataT
T DataT
Definition: types.hpp:184
sycl::_V1::detail::LessEqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:221
sycl::_V1::detail::LessThan::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:233
sycl::_V1::detail::NotEqualTo
Definition: types.hpp:208
sycl::_V1::detail::LessThan
Definition: types.hpp:232
sycl::_V1::detail::vec_helper< std::byte >::RetType
std::uint8_t RetType
Definition: types.hpp:154
sycl::_V1::detail::GetScalarOp::DataT
T DataT
Definition: types.hpp:194
sycl::_V1::ext::intel::experimental::byte
unsigned char byte
Definition: online_compiler.hpp:22
sycl::_V1::detail::vec_helper
Definition: types.hpp:142
type_list.hpp