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