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 
13 #include <CL/sycl/detail/generic_type_traits.hpp>
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 <CL/sycl/access/access.hpp>
50 #include <CL/sycl/aliases.hpp>
51 #include <CL/sycl/detail/common.hpp>
52 #include <CL/sycl/detail/helpers.hpp>
53 #include <CL/sycl/detail/type_traits.hpp>
54 #include <CL/sycl/half_type.hpp>
55 #include <CL/sycl/marray.hpp>
56 #include <CL/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 __SYCL_INLINE_NAMESPACE(cl) {
69 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>
216 using is_int_to_int =
217  std::integral_constant<bool, std::is_integral<T>::value &&
218  std::is_integral<R>::value>;
219 
220 template <typename T, typename R>
221 using is_sint_to_sint =
222  std::integral_constant<bool, is_sigeninteger<T>::value &&
223  is_sigeninteger<R>::value>;
224 
225 template <typename T, typename R>
226 using is_uint_to_uint =
227  std::integral_constant<bool, is_sugeninteger<T>::value &&
228  is_sugeninteger<R>::value>;
229 
230 template <typename T, typename R>
231 using is_sint_to_from_uint = std::integral_constant<
232  bool, (is_sugeninteger<T>::value && is_sigeninteger<R>::value) ||
233  (is_sigeninteger<T>::value && is_sugeninteger<R>::value)>;
234 
235 template <typename T, typename R>
236 using is_sint_to_float =
237  std::integral_constant<bool, std::is_integral<T>::value &&
238  !(std::is_unsigned<T>::value) &&
239  detail::is_floating_point<R>::value>;
240 
241 template <typename T, typename R>
242 using is_uint_to_float =
243  std::integral_constant<bool, std::is_unsigned<T>::value &&
244  detail::is_floating_point<R>::value>;
245 
246 template <typename T, typename R>
247 using is_int_to_float =
248  std::integral_constant<bool, std::is_integral<T>::value &&
249  detail::is_floating_point<R>::value>;
250 
251 template <typename T, typename R>
252 using is_float_to_int =
253  std::integral_constant<bool, detail::is_floating_point<T>::value &&
254  std::is_integral<R>::value>;
255 
256 template <typename T, typename R>
257 using is_float_to_float =
258  std::integral_constant<bool, detail::is_floating_point<T>::value &&
259  detail::is_floating_point<R>::value>;
260 template <typename T>
261 using is_standard_type =
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 detail::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 detail::enable_if_t<!std::is_same<T, R>::value &&
277  (is_int_to_int<T, R>::value ||
278  is_int_to_float<T, R>::value ||
279  is_float_to_float<T, R>::value),
280  R>
281 convertImpl(T Value) {
282  return static_cast<R>(Value);
283 }
284 
285 // float to int
286 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
287  typename OpenCLR>
288 detail::enable_if_t<is_float_to_int<T, R>::value, R> convertImpl(T Value) {
289  switch (roundingMode) {
290  // Round to nearest even is default rounding mode for floating-point types
291  case rounding_mode::automatic:
292  // Round to nearest even.
293  case rounding_mode::rte: {
294  int OldRoundingDirection = std::fegetround();
295  int Err = std::fesetround(FE_TONEAREST);
296  if (Err)
297  throw runtime_error("Unable to set rounding mode to FE_TONEAREST",
298  PI_ERROR_UNKNOWN);
299  R Result = std::rint(Value);
300  Err = std::fesetround(OldRoundingDirection);
301  if (Err)
302  throw runtime_error("Unable to restore rounding mode.", PI_ERROR_UNKNOWN);
303  return Result;
304  }
305  // Round toward zero.
306  case rounding_mode::rtz:
307  return std::trunc(Value);
308  // Round toward positive infinity.
309  case rounding_mode::rtp:
310  return std::ceil(Value);
311  // Round toward negative infinity.
312  case rounding_mode::rtn:
313  return std::floor(Value);
314  };
315  assert(false && "Unsupported rounding mode!");
316  return static_cast<R>(Value);
317 }
318 #else
319 
320 template <rounding_mode Mode>
321 using RteOrAutomatic = detail::bool_constant<Mode == rounding_mode::automatic ||
322  Mode == rounding_mode::rte>;
323 
324 template <rounding_mode Mode>
325 using Rtz = detail::bool_constant<Mode == rounding_mode::rtz>;
326 
327 template <rounding_mode Mode>
328 using Rtp = detail::bool_constant<Mode == rounding_mode::rtp>;
329 
330 template <rounding_mode Mode>
331 using Rtn = detail::bool_constant<Mode == rounding_mode::rtn>;
332 
333 // convert types with an equal size and diff names
334 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
335  typename OpenCLR>
336 detail::enable_if_t<
337  !std::is_same<T, R>::value && std::is_same<OpenCLT, OpenCLR>::value, R>
338 convertImpl(T Value) {
339  return static_cast<R>(Value);
340 }
341 
342 // signed to signed
343 #define __SYCL_GENERATE_CONVERT_IMPL(DestType) \
344  template <typename T, typename R, rounding_mode roundingMode, \
345  typename OpenCLT, typename OpenCLR> \
346  detail::enable_if_t<is_sint_to_sint<T, R>::value && \
347  !std::is_same<OpenCLT, OpenCLR>::value && \
348  (std::is_same<OpenCLR, cl_##DestType>::value || \
349  (std::is_same<OpenCLR, signed char>::value && \
350  std::is_same<DestType, char>::value)), \
351  R> \
352  convertImpl(T Value) { \
353  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
354  return __spirv_SConvert##_R##DestType(OpValue); \
355  }
356 
357 __SYCL_GENERATE_CONVERT_IMPL(char)
358 __SYCL_GENERATE_CONVERT_IMPL(short)
359 __SYCL_GENERATE_CONVERT_IMPL(int)
360 __SYCL_GENERATE_CONVERT_IMPL(long)
361 
362 #undef __SYCL_GENERATE_CONVERT_IMPL
363 
364 // unsigned to unsigned
365 #define __SYCL_GENERATE_CONVERT_IMPL(DestType) \
366  template <typename T, typename R, rounding_mode roundingMode, \
367  typename OpenCLT, typename OpenCLR> \
368  detail::enable_if_t<is_uint_to_uint<T, R>::value && \
369  !std::is_same<OpenCLT, OpenCLR>::value && \
370  std::is_same<OpenCLR, cl_##DestType>::value, \
371  R> \
372  convertImpl(T Value) { \
373  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
374  return __spirv_UConvert##_R##DestType(OpValue); \
375  }
376 
377 __SYCL_GENERATE_CONVERT_IMPL(uchar)
378 __SYCL_GENERATE_CONVERT_IMPL(ushort)
379 __SYCL_GENERATE_CONVERT_IMPL(uint)
380 __SYCL_GENERATE_CONVERT_IMPL(ulong)
381 
382 #undef __SYCL_GENERATE_CONVERT_IMPL
383 
384 // unsigned to (from) signed
385 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
386  typename OpenCLR>
387 detail::enable_if_t<is_sint_to_from_uint<T, R>::value &&
388  is_standard_type<OpenCLT>::value &&
389  is_standard_type<OpenCLR>::value,
390  R>
391 convertImpl(T Value) {
392  return static_cast<R>(Value);
393 }
394 
395 // sint to float
396 #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
397  template <typename T, typename R, rounding_mode roundingMode, \
398  typename OpenCLT, typename OpenCLR> \
399  detail::enable_if_t<is_sint_to_float<T, R>::value && \
400  (std::is_same<OpenCLR, DestType>::value || \
401  (std::is_same<OpenCLR, _Float16>::value && \
402  std::is_same<DestType, half>::value)), \
403  R> \
404  convertImpl(T Value) { \
405  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
406  return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
407  }
408 
409 __SYCL_GENERATE_CONVERT_IMPL(SToF, half)
410 __SYCL_GENERATE_CONVERT_IMPL(SToF, float)
411 __SYCL_GENERATE_CONVERT_IMPL(SToF, double)
412 
413 #undef __SYCL_GENERATE_CONVERT_IMPL
414 
415 // uint to float
416 #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
417  template <typename T, typename R, rounding_mode roundingMode, \
418  typename OpenCLT, typename OpenCLR> \
419  detail::enable_if_t<is_uint_to_float<T, R>::value && \
420  (std::is_same<OpenCLR, DestType>::value || \
421  (std::is_same<OpenCLR, _Float16>::value && \
422  std::is_same<DestType, half>::value)), \
423  R> \
424  convertImpl(T Value) { \
425  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
426  return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
427  }
428 
429 __SYCL_GENERATE_CONVERT_IMPL(UToF, half)
430 __SYCL_GENERATE_CONVERT_IMPL(UToF, float)
431 __SYCL_GENERATE_CONVERT_IMPL(UToF, double)
432 
433 #undef __SYCL_GENERATE_CONVERT_IMPL
434 
435 // float to float
436 #define __SYCL_GENERATE_CONVERT_IMPL(DestType, RoundingMode, \
437  RoundingModeCondition) \
438  template <typename T, typename R, rounding_mode roundingMode, \
439  typename OpenCLT, typename OpenCLR> \
440  detail::enable_if_t<is_float_to_float<T, R>::value && \
441  !std::is_same<OpenCLT, OpenCLR>::value && \
442  (std::is_same<OpenCLR, DestType>::value || \
443  (std::is_same<OpenCLR, _Float16>::value && \
444  std::is_same<DestType, half>::value)) && \
445  RoundingModeCondition<roundingMode>::value, \
446  R> \
447  convertImpl(T Value) { \
448  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
449  return __spirv_FConvert##_R##DestType##_##RoundingMode(OpValue); \
450  }
451 
452 #define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \
453  RoundingModeCondition) \
454  __SYCL_GENERATE_CONVERT_IMPL(double, RoundingMode, RoundingModeCondition) \
455  __SYCL_GENERATE_CONVERT_IMPL(float, RoundingMode, RoundingModeCondition) \
456  __SYCL_GENERATE_CONVERT_IMPL(half, RoundingMode, RoundingModeCondition)
457 
458 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic)
459 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz)
460 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp)
461 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)
462 
463 #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
464 #undef __SYCL_GENERATE_CONVERT_IMPL
465 
466 // float to int
467 #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \
468  RoundingModeCondition) \
469  template <typename T, typename R, rounding_mode roundingMode, \
470  typename OpenCLT, typename OpenCLR> \
471  detail::enable_if_t<is_float_to_int<T, R>::value && \
472  (std::is_same<OpenCLR, cl_##DestType>::value || \
473  (std::is_same<OpenCLR, signed char>::value && \
474  std::is_same<DestType, char>::value)) && \
475  RoundingModeCondition<roundingMode>::value, \
476  R> \
477  convertImpl(T Value) { \
478  OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
479  return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue); \
480  }
481 
482 #define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \
483  RoundingModeCondition) \
484  __SYCL_GENERATE_CONVERT_IMPL(FToS, int, RoundingMode, RoundingModeCondition) \
485  __SYCL_GENERATE_CONVERT_IMPL(FToS, char, RoundingMode, \
486  RoundingModeCondition) \
487  __SYCL_GENERATE_CONVERT_IMPL(FToS, short, RoundingMode, \
488  RoundingModeCondition) \
489  __SYCL_GENERATE_CONVERT_IMPL(FToS, long, RoundingMode, \
490  RoundingModeCondition) \
491  __SYCL_GENERATE_CONVERT_IMPL(FToU, uint, RoundingMode, \
492  RoundingModeCondition) \
493  __SYCL_GENERATE_CONVERT_IMPL(FToU, uchar, RoundingMode, \
494  RoundingModeCondition) \
495  __SYCL_GENERATE_CONVERT_IMPL(FToU, ushort, RoundingMode, \
496  RoundingModeCondition) \
497  __SYCL_GENERATE_CONVERT_IMPL(FToU, ulong, RoundingMode, RoundingModeCondition)
498 
499 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic)
500 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz)
501 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp)
502 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)
503 
504 #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
505 #undef __SYCL_GENERATE_CONVERT_IMPL
506 
507 // Back up
508 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
509  typename OpenCLR>
510 detail::enable_if_t<
511  ((!is_standard_type<T>::value && !is_standard_type<OpenCLT>::value) ||
512  (!is_standard_type<R>::value && !is_standard_type<OpenCLR>::value)) &&
513  !std::is_same<OpenCLT, OpenCLR>::value,
514  R>
515 convertImpl(T Value) {
516  return static_cast<R>(Value);
517 }
518 
519 #endif // __SYCL_DEVICE_ONLY__
520 
521 // Forward declarations
522 template <typename TransformedArgType, int Dims, typename KernelType>
523 class RoundedRangeKernel;
524 template <typename TransformedArgType, int Dims, typename KernelType>
525 class RoundedRangeKernelWithKH;
526 
527 } // namespace detail
528 
529 template <typename T> using vec_data = detail::vec_helper<T>;
530 
531 template <typename T>
532 using vec_data_t = typename detail::vec_helper<T>::RetType;
533 
534 #if defined(_WIN32) && (_MSC_VER)
535 // MSVC Compiler doesn't allow using of function arguments with alignment
536 // requirements. MSVC Compiler Error C2719: 'parameter': formal parameter with
537 // __declspec(align('#')) won't be aligned. The align __declspec modifier
538 // is not permitted on function parameters. Function parameter alignment
539 // is controlled by the calling convention used.
540 // For more information, see Calling Conventions
541 // (https://docs.microsoft.com/en-us/cpp/cpp/calling-conventions).
542 // For information on calling conventions for x64 processors, see
543 // Calling Convention
544 // (https://docs.microsoft.com/en-us/cpp/build/x64-calling-convention).
545 #pragma message("Alignment of class vec is not in accordance with SYCL \
546 specification requirements, a limitation of the MSVC compiler(Error C2719).\
547 Requested alignment applied, limited at 64.")
548 #define __SYCL_ALIGNED_VAR(type, x, var) \
549  type __declspec(align((x < 64) ? x : 64)) var
550 #else
551 #define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
552 #endif
553 
558 template <typename Type, int NumElements> class vec {
559  using DataT = Type;
560 
561  // This represent type of underlying value. There should be only one field
562  // in the class, so vec<float, 16> should be equal to float16 in memory.
563  using DataType =
564  typename detail::BaseCLTypeConverter<DataT, NumElements>::DataType;
565 
566  static constexpr int getNumElements() { return NumElements; }
567 
568  // SizeChecker is needed for vec(const argTN &... args) ctor to validate args.
569  template <int Counter, int MaxValue, class...>
570  struct SizeChecker: detail::conditional_t<Counter == MaxValue,
571  std::true_type, std::false_type> {};
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 detail::enable_if_t<
660  conjunction<TypeChecker<argTN, DataT>...>::value>;
661 
662  template <typename... argTN>
663  using EnableIfSuitableNumElements = typename detail::enable_if_t<
664  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 detail::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 detail::enable_if_t<
702  !std::is_same<DataT, cl::sycl::detail::half_impl::half>::value ||
703  !std::is_same<cl::sycl::detail::half_impl::StorageT,
704  cl::sycl::detail::host_half_impl::half_v2>::value,
705  T>;
706  template <typename T = void>
707  using EnableIfHostHalf = typename detail::enable_if_t<
708  std::is_same<DataT, cl::sycl::detail::half_impl::half>::value &&
709  std::is_same<cl::sycl::detail::half_impl::StorageT,
710  cl::sycl::detail::host_half_impl::half_v2>::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 detail::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 detail::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 detail::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 detail::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 detail::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 detail::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>
925  struct Indexer {
926  static constexpr int value = Index;
927  };
928 
929 public:
930 #ifdef __SYCL_ACCESS_RETURN
931 #error "Undefine __SYCL_ACCESS_RETURN macro"
932 #endif
933 #define __SYCL_ACCESS_RETURN this
934 #include "swizzles.def"
935 #undef __SYCL_ACCESS_RETURN
936  // End of hi/lo, even/odd, xyzw, and rgba swizzles.
937 
938  template <access::address_space Space>
939  void load(size_t Offset, multi_ptr<const DataT, Space> Ptr) {
940  for (int I = 0; I < NumElements; I++) {
941  setValue(I,
942  *multi_ptr<const DataT, Space>(Ptr + Offset * NumElements + I));
943  }
944  }
945  template <access::address_space Space>
946  void load(size_t Offset, multi_ptr<DataT, Space> Ptr) {
947  multi_ptr<const DataT, Space> ConstPtr(Ptr);
948  load(Offset, ConstPtr);
949  }
950  template <int Dimensions, access::mode Mode,
951  access::placeholder IsPlaceholder, access::target Target,
952  typename PropertyListT>
953  void
954  load(size_t Offset,
955  accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
956  Acc) {
957  multi_ptr<const DataT, detail::TargetToAS<Target>::AS> MultiPtr(Acc);
958  load(Offset, MultiPtr);
959  }
960  template <access::address_space Space>
961  void store(size_t Offset, multi_ptr<DataT, Space> Ptr) const {
962  for (int I = 0; I < NumElements; I++) {
963  *multi_ptr<DataT, Space>(Ptr + Offset * NumElements + I) = getValue(I);
964  }
965  }
966  template <int Dimensions, access::mode Mode,
967  access::placeholder IsPlaceholder, access::target Target,
968  typename PropertyListT>
969  void
970  store(size_t Offset,
971  accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
972  Acc) {
973  multi_ptr<DataT, detail::TargetToAS<Target>::AS> MultiPtr(Acc);
974  store(Offset, MultiPtr);
975  }
976 
977 #ifdef __SYCL_BINOP
978 #error "Undefine __SYCL_BINOP macro"
979 #endif
980 
981 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
982 #define __SYCL_BINOP(BINOP, OPASSIGN) \
983  template <typename Ty = vec> \
984  vec operator BINOP(const EnableIfNotHostHalf<Ty> &Rhs) const { \
985  vec Ret; \
986  Ret.m_Data = m_Data BINOP Rhs.m_Data; \
987  return Ret; \
988  } \
989  template <typename Ty = vec> \
990  vec operator BINOP(const EnableIfHostHalf<Ty> &Rhs) const { \
991  vec Ret; \
992  for (size_t I = 0; I < NumElements; ++I) { \
993  Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
994  } \
995  return Ret; \
996  } \
997  template <typename T> \
998  typename detail::enable_if_t< \
999  std::is_convertible<DataT, T>::value && \
1000  (std::is_fundamental<vec_data_t<T>>::value || \
1001  std::is_same<typename detail::remove_const_t<T>, half>::value), \
1002  vec> \
1003  operator BINOP(const T &Rhs) const { \
1004  return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1005  } \
1006  vec &operator OPASSIGN(const vec &Rhs) { \
1007  *this = *this BINOP Rhs; \
1008  return *this; \
1009  } \
1010  template <int Num = NumElements> \
1011  typename detail::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1012  const DataT &Rhs) { \
1013  *this = *this BINOP vec(Rhs); \
1014  return *this; \
1015  }
1016 #else // __SYCL_USE_EXT_VECTOR_TYPE__
1017 #define __SYCL_BINOP(BINOP, OPASSIGN) \
1018  vec operator BINOP(const vec &Rhs) const { \
1019  vec Ret; \
1020  for (size_t I = 0; I < NumElements; ++I) { \
1021  Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
1022  } \
1023  return Ret; \
1024  } \
1025  template <typename T> \
1026  typename detail::enable_if_t< \
1027  std::is_convertible<DataT, T>::value && \
1028  (std::is_fundamental<vec_data_t<T>>::value || \
1029  std::is_same<typename detail::remove_const_t<T>, half>::value), \
1030  vec> \
1031  operator BINOP(const T &Rhs) const { \
1032  return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1033  } \
1034  vec &operator OPASSIGN(const vec &Rhs) { \
1035  *this = *this BINOP Rhs; \
1036  return *this; \
1037  } \
1038  template <int Num = NumElements> \
1039  typename detail::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1040  const DataT &Rhs) { \
1041  *this = *this BINOP vec(Rhs); \
1042  return *this; \
1043  }
1044 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
1045 
1046  __SYCL_BINOP(+, +=)
1047  __SYCL_BINOP(-, -=)
1048  __SYCL_BINOP(*, *=)
1049  __SYCL_BINOP(/, /=)
1050 
1051  // TODO: The following OPs are available only when: DataT != cl_float &&
1052  // DataT != cl_double && DataT != cl_half
1053  __SYCL_BINOP(%, %=)
1054  __SYCL_BINOP(|, |=)
1055  __SYCL_BINOP(&, &=)
1056  __SYCL_BINOP(^, ^=)
1057  __SYCL_BINOP(>>, >>=)
1058  __SYCL_BINOP(<<, <<=)
1059 #undef __SYCL_BINOP
1060 #undef __SYCL_BINOP_HELP
1061 
1062  // Note: vec<>/SwizzleOp logical value is 0/-1 logic, as opposed to 0/1 logic.
1063  // As far as CTS validation is concerned, 0/-1 logic also applies when
1064  // NumElements is equal to one, which is somewhat inconsistent with being
1065  // transparent with scalar data.
1066  // TODO: Determine if vec<, NumElements=1> is needed at all, remove this
1067  // inconsistency if not by disallowing one-element vectors (as in OpenCL)
1068 
1069 #ifdef __SYCL_RELLOGOP
1070 #error "Undefine __SYCL_RELLOGOP macro"
1071 #endif
1072 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1073 // by SYCL device compiler only.
1074 #ifdef __SYCL_DEVICE_ONLY__
1075 #define __SYCL_RELLOGOP(RELLOGOP) \
1076  vec<rel_t, NumElements> operator RELLOGOP(const vec &Rhs) const { \
1077  auto Ret = \
1078  vec<rel_t, NumElements>((typename vec<rel_t, NumElements>::vector_t)( \
1079  m_Data RELLOGOP Rhs.m_Data)); \
1080  if (NumElements == 1) /*Scalar 0/1 logic was applied, invert*/ \
1081  Ret *= -1; \
1082  return Ret; \
1083  } \
1084  template <typename T> \
1085  typename detail::enable_if_t< \
1086  std::is_convertible<T, DataT>::value && \
1087  (std::is_fundamental<vec_data_t<T>>::value || \
1088  std::is_same<T, half>::value), \
1089  vec<rel_t, NumElements>> \
1090  operator RELLOGOP(const T &Rhs) const { \
1091  return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1092  }
1093 #else
1094 #define __SYCL_RELLOGOP(RELLOGOP) \
1095  vec<rel_t, NumElements> operator RELLOGOP(const vec &Rhs) const { \
1096  vec<rel_t, NumElements> Ret; \
1097  for (size_t I = 0; I < NumElements; ++I) { \
1098  Ret.setValue(I, -(vec_data<DataT>::get(getValue(I)) \
1099  RELLOGOP vec_data<DataT>::get(Rhs.getValue(I)))); \
1100  } \
1101  return Ret; \
1102  } \
1103  template <typename T> \
1104  typename detail::enable_if_t< \
1105  std::is_convertible<T, DataT>::value && \
1106  (std::is_fundamental<vec_data_t<T>>::value || \
1107  std::is_same<T, half>::value), \
1108  vec<rel_t, NumElements>> \
1109  operator RELLOGOP(const T &Rhs) const { \
1110  return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1111  }
1112 #endif
1113 
1114  __SYCL_RELLOGOP(==)
1115  __SYCL_RELLOGOP(!=)
1116  __SYCL_RELLOGOP(>)
1117  __SYCL_RELLOGOP(<)
1118  __SYCL_RELLOGOP(>=)
1119  __SYCL_RELLOGOP(<=)
1120  // TODO: limit to integral types.
1121  __SYCL_RELLOGOP(&&)
1122  __SYCL_RELLOGOP(||)
1123 #undef __SYCL_RELLOGOP
1124 
1125 #ifdef __SYCL_UOP
1126 #error "Undefine __SYCL_UOP macro"
1127 #endif
1128 #define __SYCL_UOP(UOP, OPASSIGN) \
1129  vec &operator UOP() { \
1130  *this OPASSIGN vec_data<DataT>::get(1); \
1131  return *this; \
1132  } \
1133  vec operator UOP(int) { \
1134  vec Ret(*this); \
1135  *this OPASSIGN vec_data<DataT>::get(1); \
1136  return Ret; \
1137  }
1138 
1139  __SYCL_UOP(++, +=)
1140  __SYCL_UOP(--, -=)
1141 #undef __SYCL_UOP
1142 
1143  // Available only when: dataT != cl_float && dataT != cl_double
1144  // && dataT != cl_half
1145  template <typename T = DataT>
1146  typename detail::enable_if_t<std::is_integral<vec_data_t<T>>::value, vec>
1147  operator~() const {
1148 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1149 // by SYCL device compiler only.
1150 #ifdef __SYCL_DEVICE_ONLY__
1151  return vec{
1152  (typename vec::DataType)~m_Data};
1153 #else
1154  vec Ret;
1155  for (size_t I = 0; I < NumElements; ++I) {
1156  Ret.setValue(I, ~getValue(I));
1157  }
1158  return Ret;
1159 #endif
1160  }
1161 
1162  vec<rel_t, NumElements> operator!() const {
1163 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1164 // by SYCL device compiler only.
1165 #ifdef __SYCL_DEVICE_ONLY__
1166  return vec<rel_t, NumElements>{
1167  (typename vec<rel_t, NumElements>::DataType)!m_Data};
1168 #else
1169  vec<rel_t, NumElements> Ret;
1170  for (size_t I = 0; I < NumElements; ++I) {
1171  Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1172  }
1173  return Ret;
1174 #endif
1175  }
1176 
1177  vec operator+() const {
1178 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1179 // by SYCL device compiler only.
1180 #ifdef __SYCL_DEVICE_ONLY__
1181  return vec{+m_Data};
1182 #else
1183  vec Ret;
1184  for (size_t I = 0; I < NumElements; ++I)
1185  Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1186  return Ret;
1187 #endif
1188  }
1189 
1190  vec operator-() const {
1191 // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined
1192 // by SYCL device compiler only.
1193 #ifdef __SYCL_DEVICE_ONLY__
1194  return vec{-m_Data};
1195 #else
1196  vec Ret;
1197  for (size_t I = 0; I < NumElements; ++I)
1198  Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
1199  return Ret;
1200 #endif
1201  }
1202 
1203  // OP is: &&, ||
1204  // vec<RET, NumElements> operatorOP(const vec<DataT, NumElements> &Rhs) const;
1205  // vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
1206 
1207  // OP is: ==, !=, <, >, <=, >=
1208  // vec<RET, NumElements> operatorOP(const vec<DataT, NumElements> &Rhs) const;
1209  // vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
1210 private:
1211  // Generic method that execute "Operation" on underlying values.
1212 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1213  template <template <typename> class Operation,
1214  typename Ty = vec<DataT, NumElements>>
1215  vec<DataT, NumElements>
1216  operatorHelper(const EnableIfNotHostHalf<Ty> &Rhs) const {
1217  vec<DataT, NumElements> Result;
1218  Operation<DataType> Op;
1219  Result.m_Data = Op(m_Data, Rhs.m_Data);
1220  return Result;
1221  }
1222 
1223  template <template <typename> class Operation,
1224  typename Ty = vec<DataT, NumElements>>
1225  vec<DataT, NumElements>
1226  operatorHelper(const EnableIfHostHalf<Ty> &Rhs) const {
1227  vec<DataT, NumElements> Result;
1228  Operation<DataT> Op;
1229  for (size_t I = 0; I < NumElements; ++I) {
1230  Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1231  }
1232  return Result;
1233  }
1234 #else // __SYCL_USE_EXT_VECTOR_TYPE__
1235  template <template <typename> class Operation>
1236  vec<DataT, NumElements>
1237  operatorHelper(const vec<DataT, NumElements> &Rhs) const {
1238  vec<DataT, NumElements> Result;
1239  Operation<DataT> Op;
1240  for (size_t I = 0; I < NumElements; ++I) {
1241  Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1242  }
1243  return Result;
1244  }
1245 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
1246 
1247 // setValue and getValue should be able to operate on different underlying
1248 // types: enum cl_float#N , builtin vector float#N, builtin type float.
1249 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1250  template <int Num = NumElements, typename Ty = int,
1251  typename = typename detail::enable_if_t<1 != Num>>
1252  constexpr void setValue(EnableIfNotHostHalf<Ty> Index, const DataT &Value,
1253  int) {
1254  m_Data[Index] = vec_data<DataT>::get(Value);
1255  }
1256 
1257  template <int Num = NumElements, typename Ty = int,
1258  typename = typename detail::enable_if_t<1 != Num>>
1259  DataT getValue(EnableIfNotHostHalf<Ty> Index, int) const {
1260  return vec_data<DataT>::get(m_Data[Index]);
1261  }
1262 
1263  template <int Num = NumElements, typename Ty = int,
1264  typename = typename detail::enable_if_t<1 != Num>>
1265  constexpr void setValue(EnableIfHostHalf<Ty> Index, const DataT &Value, int) {
1266  m_Data.s[Index] = vec_data<DataT>::get(Value);
1267  }
1268 
1269  template <int Num = NumElements, typename Ty = int,
1270  typename = typename detail::enable_if_t<1 != Num>>
1271  DataT getValue(EnableIfHostHalf<Ty> Index, int) const {
1272  return vec_data<DataT>::get(m_Data.s[Index]);
1273  }
1274 #else // __SYCL_USE_EXT_VECTOR_TYPE__
1275  template <int Num = NumElements,
1276  typename = typename detail::enable_if_t<1 != Num>>
1277  constexpr void setValue(int Index, const DataT &Value, int) {
1278  m_Data.s[Index] = vec_data<DataT>::get(Value);
1279  }
1280 
1281  template <int Num = NumElements,
1282  typename = typename detail::enable_if_t<1 != Num>>
1283  DataT getValue(int Index, int) const {
1284  return vec_data<DataT>::get(m_Data.s[Index]);
1285  }
1286 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
1287 
1288  template <int Num = NumElements,
1289  typename = typename detail::enable_if_t<1 == Num>>
1290  constexpr void setValue(int, const DataT &Value, float) {
1291  m_Data = vec_data<DataT>::get(Value);
1292  }
1293 
1294  template <int Num = NumElements,
1295  typename = typename detail::enable_if_t<1 == Num>>
1296  DataT getValue(int, float) const {
1297  return vec_data<DataT>::get(m_Data);
1298  }
1299 
1300  // Special proxies as specialization is not allowed in class scope.
1301  constexpr void setValue(int Index, const DataT &Value) {
1302  if (NumElements == 1)
1303  setValue(Index, Value, 0);
1304  else
1305  setValue(Index, Value, 0.f);
1306  }
1307 
1308  DataT getValue(int Index) const {
1309  return (NumElements == 1) ? getValue(Index, 0) : getValue(Index, 0.f);
1310  }
1311 
1312  // Helpers for variadic template constructor of vec.
1313  template <typename T, typename... argTN>
1314  constexpr int vaargCtorHelper(int Idx, const T &arg) {
1315  setValue(Idx, arg);
1316  return Idx + 1;
1317  }
1318 
1319  template <typename DataT_, int NumElements_>
1320  constexpr int vaargCtorHelper(int Idx, const vec<DataT_, NumElements_> &arg) {
1321  for (size_t I = 0; I < NumElements_; ++I) {
1322  setValue(Idx + I, arg.getValue(I));
1323  }
1324  return Idx + NumElements_;
1325  }
1326 
1327  template <typename DataT_, int NumElements_, typename T2, typename T3,
1328  template <typename> class T4, int... T5>
1329  constexpr int
1330  vaargCtorHelper(int Idx, const detail::SwizzleOp<vec<DataT_, NumElements_>,
1331  T2, T3, T4, T5...> &arg) {
1332  size_t NumElems = sizeof...(T5);
1333  for (size_t I = 0; I < NumElems; ++I) {
1334  setValue(Idx + I, arg.getValue(I));
1335  }
1336  return Idx + NumElems;
1337  }
1338 
1339  template <typename DataT_, int NumElements_, typename T2, typename T3,
1340  template <typename> class T4, int... T5>
1341  constexpr int
1342  vaargCtorHelper(int Idx,
1343  const detail::SwizzleOp<const vec<DataT_, NumElements_>, T2,
1344  T3, T4, T5...> &arg) {
1345  size_t NumElems = sizeof...(T5);
1346  for (size_t I = 0; I < NumElems; ++I) {
1347  setValue(Idx + I, arg.getValue(I));
1348  }
1349  return Idx + NumElems;
1350  }
1351 
1352  template <typename T1, typename... argTN>
1353  constexpr void vaargCtorHelper(int Idx, const T1 &arg,
1354  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 = detail::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 detail::enable_if_t<
1430  1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1431 
1432  template <int IdxNum, typename T = void>
1433  using EnableIfMultipleIndexes = typename detail::enable_if_t<
1434  1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1435 
1436  template <typename T>
1437  using EnableIfScalarType = typename detail::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 detail::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 detail::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 <typename T1, typename T2, typename T3, template <typename> class T4,
1719  int... T5,
1720  typename =
1721  typename detail::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 <typename T1, typename T2, typename T3, template <typename> class T4,
1731  int... T5,
1732  typename =
1733  typename detail::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>
1860  struct Indexer {
1861  static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...};
1862  static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index];
1863  };
1864 
1865 public:
1866 #ifdef __SYCL_ACCESS_RETURN
1867 #error "Undefine __SYCL_ACCESS_RETURN macro"
1868 #endif
1869 #define __SYCL_ACCESS_RETURN m_Vector
1870 #include "swizzles.def"
1871 #undef __SYCL_ACCESS_RETURN
1872  // End of hi/lo, even/odd, xyzw, and rgba swizzles.
1873 
1874  // Leave store() interface to automatic conversion to vec<>.
1875  // Load to vec_t and then assign to swizzle.
1876  template <access::address_space Space>
1877  void load(size_t offset, multi_ptr<DataT, Space> ptr) {
1878  vec_t Tmp;
1879  Tmp.template load(offset, ptr);
1880  *this = Tmp;
1881  }
1882 
1883  template <typename convertT, rounding_mode roundingMode>
1884  vec<convertT, sizeof...(Indexes)> convert() const {
1885  // First materialize the swizzle to vec_t and then apply convert() to it.
1886  vec_t Tmp = *this;
1887  return Tmp.template convert<convertT, roundingMode>();
1888  }
1889 
1890  template <typename asT> asT as() const {
1891  // First materialize the swizzle to vec_t and then apply as() to it.
1892  vec_t Tmp = *this;
1893  static_assert((sizeof(Tmp) == sizeof(asT)),
1894  "The new SYCL vec type must have the same storage size in "
1895  "bytes as this SYCL swizzled vec");
1896  static_assert(
1897  detail::is_contained<asT, detail::gtl::vector_basic_list>::value,
1898  "asT must be SYCL vec of a different element type and "
1899  "number of elements specified by asT");
1900  return Tmp.template as<asT>();
1901  }
1902 
1903 private:
1904  SwizzleOp(const SwizzleOp &Rhs)
1905  : m_Vector(Rhs.m_Vector), m_LeftOperation(Rhs.m_LeftOperation),
1906  m_RightOperation(Rhs.m_RightOperation) {}
1907 
1908  SwizzleOp(VecT *Vector, OperationLeftT LeftOperation,
1909  OperationRightT RightOperation)
1910  : m_Vector(Vector), m_LeftOperation(LeftOperation),
1911  m_RightOperation(RightOperation) {}
1912 
1913  SwizzleOp(VecT *Vector) : m_Vector(Vector) {}
1914 
1915  SwizzleOp(SwizzleOp &&Rhs)
1916  : m_Vector(Rhs.m_Vector), m_LeftOperation(std::move(Rhs.m_LeftOperation)),
1917  m_RightOperation(std::move(Rhs.m_RightOperation)) {}
1918 
1919  // Either performing CurrentOperation on results of left and right operands
1920  // or reading values from actual vector. Perform implicit type conversion when
1921  // the number of elements == 1
1922 
1923  template <int IdxNum = getNumElements()>
1924  CommonDataT getValue(EnableIfOneIndex<IdxNum, size_t> Index) const {
1925  if (std::is_same<OperationCurrentT<DataT>, GetOp<DataT>>::value) {
1926  std::array<int, getNumElements()> Idxs{Indexes...};
1927  return m_Vector->getValue(Idxs[Index]);
1928  }
1929  auto Op = OperationCurrentT<vec_data_t<CommonDataT>>();
1930  return vec_data<CommonDataT>::get(
1931  Op(vec_data<CommonDataT>::get(m_LeftOperation.getValue(Index)),
1932  vec_data<CommonDataT>::get(m_RightOperation.getValue(Index))));
1933  }
1934 
1935  template <int IdxNum = getNumElements()>
1936  DataT getValue(EnableIfMultipleIndexes<IdxNum, size_t> Index) const {
1937  if (std::is_same<OperationCurrentT<DataT>, GetOp<DataT>>::value) {
1938  std::array<int, getNumElements()> Idxs{Indexes...};
1939  return m_Vector->getValue(Idxs[Index]);
1940  }
1941  auto Op = OperationCurrentT<vec_data_t<DataT>>();
1942  return vec_data<DataT>::get(
1943  Op(vec_data<DataT>::get(m_LeftOperation.getValue(Index)),
1944  vec_data<DataT>::get(m_RightOperation.getValue(Index))));
1945  }
1946 
1947  template <template <typename> class Operation, typename RhsOperation>
1948  void operatorHelper(const RhsOperation &Rhs) {
1949  Operation<vec_data_t<DataT>> Op;
1950  std::array<int, getNumElements()> Idxs{Indexes...};
1951  for (size_t I = 0; I < Idxs.size(); ++I) {
1952  DataT Res = vec_data<DataT>::get(
1953  Op(vec_data<DataT>::get(m_Vector->getValue(Idxs[I])),
1954  vec_data<DataT>::get(Rhs.getValue(I))));
1955  m_Vector->setValue(Idxs[I], Res);
1956  }
1957  }
1958 
1959  // fields
1960  VecT *m_Vector;
1961 
1962  OperationLeftT m_LeftOperation;
1963  OperationRightT m_RightOperation;
1964 
1965  // friends
1966  template <typename T1, int T2> friend class cl::sycl::vec;
1967 
1968  template <typename T1, typename T2, typename T3, template <typename> class T4,
1969  int... T5>
1970  friend class SwizzleOp;
1971 };
1972 } // namespace detail
1973 
1974 // scalar BINOP vec<>
1975 // scalar BINOP SwizzleOp
1976 // vec<> BINOP SwizzleOp
1977 #ifdef __SYCL_BINOP
1978 #error "Undefine __SYCL_BINOP macro"
1979 #endif
1980 #define __SYCL_BINOP(BINOP) \
1981  template <typename T, int Num> \
1982  typename detail::enable_if_t< \
1983  std::is_fundamental<vec_data_t<T>>::value || \
1984  std::is_same<typename detail::remove_const_t<T>, half>::value, \
1985  vec<T, Num>> \
1986  operator BINOP(const T &Lhs, const vec<T, Num> &Rhs) { \
1987  return vec<T, Num>(Lhs) BINOP Rhs; \
1988  } \
1989  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
1990  template <typename> class OperationCurrentT, int... Indexes, \
1991  typename T, typename T1 = typename VecT::element_type, \
1992  int Num = sizeof...(Indexes)> \
1993  typename detail::enable_if_t< \
1994  std::is_convertible<T, T1>::value && \
1995  (std::is_fundamental<vec_data_t<T>>::value || \
1996  std::is_same<typename detail::remove_const_t<T>, half>::value), \
1997  vec<T1, Num>> \
1998  operator BINOP( \
1999  const T &Lhs, \
2000  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2001  OperationCurrentT, Indexes...> &Rhs) { \
2002  vec<T1, Num> Tmp = Rhs; \
2003  return Lhs BINOP Tmp; \
2004  } \
2005  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2006  template <typename> class OperationCurrentT, int... Indexes, \
2007  typename T = typename VecT::element_type, \
2008  int Num = sizeof...(Indexes)> \
2009  vec<T, Num> operator BINOP( \
2010  const vec<T, Num> &Lhs, \
2011  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2012  OperationCurrentT, Indexes...> &Rhs) { \
2013  vec<T, Num> Tmp = Rhs; \
2014  return Lhs BINOP Tmp; \
2015  }
2016 
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 __SYCL_BINOP(<<)
2026 #undef __SYCL_BINOP
2027 
2028 // scalar RELLOGOP vec<>
2029 // scalar RELLOGOP SwizzleOp
2030 // vec<> RELLOGOP SwizzleOp
2031 #ifdef __SYCL_RELLOGOP
2032 #error "Undefine __SYCL_RELLOGOP macro"
2033 #endif
2034 #define __SYCL_RELLOGOP(RELLOGOP) \
2035  template <typename T, typename DataT, int Num> \
2036  typename detail::enable_if_t< \
2037  std::is_convertible<T, DataT>::value && \
2038  (std::is_fundamental<vec_data_t<T>>::value || \
2039  std::is_same<typename detail::remove_const_t<T>, half>::value), \
2040  vec<detail::rel_t<DataT>, Num>> \
2041  operator RELLOGOP(const T &Lhs, const vec<DataT, Num> &Rhs) { \
2042  return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs; \
2043  } \
2044  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2045  template <typename> class OperationCurrentT, int... Indexes, \
2046  typename T, typename T1 = typename VecT::element_type, \
2047  int Num = sizeof...(Indexes)> \
2048  typename detail::enable_if_t< \
2049  std::is_convertible<T, T1>::value && \
2050  (std::is_fundamental<vec_data_t<T>>::value || \
2051  std::is_same<typename detail::remove_const_t<T>, half>::value), \
2052  vec<detail::rel_t<T1>, Num>> \
2053  operator RELLOGOP( \
2054  const T &Lhs, \
2055  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2056  OperationCurrentT, Indexes...> &Rhs) { \
2057  vec<T1, Num> Tmp = Rhs; \
2058  return Lhs RELLOGOP Tmp; \
2059  } \
2060  template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2061  template <typename> class OperationCurrentT, int... Indexes, \
2062  typename T = typename VecT::element_type, \
2063  int Num = sizeof...(Indexes)> \
2064  vec<detail::rel_t<T>, Num> operator RELLOGOP( \
2065  const vec<T, Num> &Lhs, \
2066  const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2067  OperationCurrentT, Indexes...> &Rhs) { \
2068  vec<T, Num> Tmp = Rhs; \
2069  return Lhs RELLOGOP Tmp; \
2070  }
2071 
2072 __SYCL_RELLOGOP(==)
2073 __SYCL_RELLOGOP(!=)
2074 __SYCL_RELLOGOP(>)
2075 __SYCL_RELLOGOP(<)
2076 __SYCL_RELLOGOP(>=)
2077 __SYCL_RELLOGOP(<=)
2078 // TODO: limit to integral types.
2079 __SYCL_RELLOGOP(&&)
2080 __SYCL_RELLOGOP(||)
2081 #undef __SYCL_RELLOGOP
2082 
2083 } // namespace sycl
2084 } // __SYCL_INLINE_NAMESPACE(cl)
2085 
2086 
2087 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
2088 #define __SYCL_DECLARE_TYPE_VIA_CL_T(type) \
2089  using __##type##_t = cl::sycl::cl_##type; \
2090  using __##type##2_vec_t = \
2091  cl::sycl::cl_##type __attribute__((ext_vector_type(2))); \
2092  using __##type##3_vec_t = \
2093  cl::sycl::cl_##type __attribute__((ext_vector_type(3))); \
2094  using __##type##4_vec_t = \
2095  cl::sycl::cl_##type __attribute__((ext_vector_type(4))); \
2096  using __##type##8_vec_t = \
2097  cl::sycl::cl_##type __attribute__((ext_vector_type(8))); \
2098  using __##type##16_vec_t = \
2099  cl::sycl::cl_##type __attribute__((ext_vector_type(16)));
2100 
2101 #define __SYCL_DECLARE_TYPE_T(type) \
2102  using __##type##_t = cl::sycl::type; \
2103  using __##type##2_vec_t = \
2104  cl::sycl::type __attribute__((ext_vector_type(2))); \
2105  using __##type##3_vec_t = \
2106  cl::sycl::type __attribute__((ext_vector_type(3))); \
2107  using __##type##4_vec_t = \
2108  cl::sycl::type __attribute__((ext_vector_type(4))); \
2109  using __##type##8_vec_t = \
2110  cl::sycl::type __attribute__((ext_vector_type(8))); \
2111  using __##type##16_vec_t = \
2112  cl::sycl::type __attribute__((ext_vector_type(16)));
2113 
2114 __SYCL_DECLARE_TYPE_VIA_CL_T(char)
2115 __SYCL_DECLARE_TYPE_T(schar)
2116 __SYCL_DECLARE_TYPE_VIA_CL_T(uchar)
2117 __SYCL_DECLARE_TYPE_VIA_CL_T(short)
2118 __SYCL_DECLARE_TYPE_VIA_CL_T(ushort)
2119 __SYCL_DECLARE_TYPE_VIA_CL_T(int)
2120 __SYCL_DECLARE_TYPE_VIA_CL_T(uint)
2121 __SYCL_DECLARE_TYPE_VIA_CL_T(long)
2122 __SYCL_DECLARE_TYPE_VIA_CL_T(ulong)
2123 __SYCL_DECLARE_TYPE_T(longlong)
2124 __SYCL_DECLARE_TYPE_T(ulonglong)
2125 // Note: halfs are not declared here, because they have different representation
2126 // between host and device, see separate handling below
2127 __SYCL_DECLARE_TYPE_VIA_CL_T(float)
2128 __SYCL_DECLARE_TYPE_VIA_CL_T(double)
2129 
2130 #define __SYCL_GET_CL_TYPE(target, num) __##target##num##_vec_t
2131 #define __SYCL_GET_SCALAR_CL_TYPE(target) target
2132 
2133 #undef __SYCL_DECLARE_TYPE_VIA_CL_T
2134 #undef __SYCL_DECLARE_TYPE_T
2135 #else // __SYCL_USE_EXT_VECTOR_TYPE__
2136 #define __SYCL_GET_CL_TYPE(target, num) ::cl_##target##num
2137 #define __SYCL_GET_SCALAR_CL_TYPE(target) ::cl_##target
2138 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
2139 
2140 using __half_t = cl::sycl::detail::half_impl::StorageT;
2141 using __half2_vec_t = cl::sycl::detail::half_impl::Vec2StorageT;
2142 using __half3_vec_t = cl::sycl::detail::half_impl::Vec3StorageT;
2143 using __half4_vec_t = cl::sycl::detail::half_impl::Vec4StorageT;
2144 using __half8_vec_t = cl::sycl::detail::half_impl::Vec8StorageT;
2145 using __half16_vec_t = cl::sycl::detail::half_impl::Vec16StorageT;
2146 #define __SYCL_GET_CL_HALF_TYPE(target, num) __##target##num##_vec_t
2147 
2148 __SYCL_INLINE_NAMESPACE(cl) {
2149 namespace sycl {
2150 namespace detail {
2151 // select_apply_cl_t selects from T8/T16/T32/T64 basing on
2152 // sizeof(IN). expected to handle scalar types in IN.
2153 template <typename T, typename T8, typename T16, typename T32, typename T64>
2154 using select_apply_cl_t =
2155  conditional_t<sizeof(T) == 1, T8,
2156  conditional_t<sizeof(T) == 2, T16,
2157  conditional_t<sizeof(T) == 4, T32, T64>>>;
2158 } // detail
2159 
2160 #define __SYCL_DECLARE_CONVERTER(base, num) \
2161  template <> class BaseCLTypeConverter<base, num> { \
2162  public: \
2163  using DataType = __SYCL_GET_CL_TYPE(base, num); \
2164  };
2165 
2166 #define __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, num) \
2167  template <> class BaseCLTypeConverter<base, num> { \
2168  public: \
2169  using DataType = detail::select_apply_cl_t< \
2170  base, __SYCL_GET_CL_TYPE(char, num), __SYCL_GET_CL_TYPE(short, num), \
2171  __SYCL_GET_CL_TYPE(int, num), __SYCL_GET_CL_TYPE(long, num)>; \
2172  };
2173 
2174 #define __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, num) \
2175  template <> class BaseCLTypeConverter<base, num> { \
2176  public: \
2177  using DataType = detail::select_apply_cl_t< \
2178  base, __SYCL_GET_CL_TYPE(uchar, num), __SYCL_GET_CL_TYPE(ushort, num), \
2179  __SYCL_GET_CL_TYPE(uint, num), __SYCL_GET_CL_TYPE(ulong, num)>; \
2180  };
2181 
2182 #define __SYCL_DECLARE_FLOAT_CONVERTER(base, num) \
2183  template <> class BaseCLTypeConverter<base, num> { \
2184  public: \
2185  using DataType = detail::select_apply_cl_t< \
2186  base, std::false_type, __SYCL_GET_CL_HALF_TYPE(half, num), \
2187  __SYCL_GET_CL_TYPE(float, num), __SYCL_GET_CL_TYPE(double, num)>; \
2188  };
2189 
2190 #define __SYCL_DECLARE_LONGLONG_CONVERTER(base, num) \
2191  template <> class BaseCLTypeConverter<base##long, num> { \
2192  public: \
2193  using DataType = __SYCL_GET_CL_TYPE(base, num); \
2194  };
2195 
2196 #define __SYCL_DECLARE_SCHAR_CONVERTER(num) \
2197  template <> class BaseCLTypeConverter<schar, num> { \
2198  public: \
2199  using DataType = detail::select_apply_cl_t< \
2200  schar, __SYCL_GET_CL_TYPE(char, num), __SYCL_GET_CL_TYPE(short, num), \
2201  __SYCL_GET_CL_TYPE(int, num), __SYCL_GET_CL_TYPE(long, num)>; \
2202  };
2203 
2204 #define __SYCL_DECLARE_BOOL_CONVERTER(num) \
2205  template <> class BaseCLTypeConverter<bool, num> { \
2206  public: \
2207  using DataType = detail::select_apply_cl_t< \
2208  bool, __SYCL_GET_CL_TYPE(char, num), __SYCL_GET_CL_TYPE(short, num), \
2209  __SYCL_GET_CL_TYPE(int, num), __SYCL_GET_CL_TYPE(long, num)>; \
2210  };
2211 
2212 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2213 #define __SYCL_DECLARE_BYTE_CONVERTER(num) \
2214  template <> class BaseCLTypeConverter<std::byte, num> { \
2215  public: \
2216  using DataType = __SYCL_GET_CL_TYPE(uchar, num); \
2217  };
2218 #endif
2219 #define __SYCL_DECLARE_HALF_CONVERTER(base, num) \
2220  template <> class BaseCLTypeConverter<base, num> { \
2221  public: \
2222  using DataType = __SYCL_GET_CL_HALF_TYPE(base, num); \
2223  };
2224 
2225 #define __SYCL_DECLARE_SCALAR_SCHAR_CONVERTER \
2226  template <> class BaseCLTypeConverter<schar, 1> { \
2227  public: \
2228  using DataType = schar; \
2229  };
2230 
2231 #define __SYCL_DECLARE_SCALAR_BOOL_CONVERTER \
2232  template <> class BaseCLTypeConverter<bool, 1> { \
2233  public: \
2234  using DataType = bool; \
2235  };
2236 
2237 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2238 #define __SYCL_DECLARE_SCALAR_BYTE_CONVERTER \
2239  template <> class BaseCLTypeConverter<std::byte, 1> { \
2240  public: \
2241  using DataType = __SYCL_GET_SCALAR_CL_TYPE(uchar); \
2242  };
2243 #endif
2244 #define __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2245  template <> class BaseCLTypeConverter<base, 1> { \
2246  public: \
2247  using DataType = __SYCL_GET_SCALAR_CL_TYPE(base); \
2248  };
2249 
2250 #define __SYCL_DECLARE_VECTOR_CONVERTERS(base) \
2251  namespace detail { \
2252  __SYCL_DECLARE_CONVERTER(base, 2) \
2253  __SYCL_DECLARE_CONVERTER(base, 3) \
2254  __SYCL_DECLARE_CONVERTER(base, 4) \
2255  __SYCL_DECLARE_CONVERTER(base, 8) \
2256  __SYCL_DECLARE_CONVERTER(base, 16) \
2257  __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2258  } // namespace detail
2259 
2260 #define __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(base) \
2261  namespace detail { \
2262  __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 2) \
2263  __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 3) \
2264  __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 4) \
2265  __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 8) \
2266  __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 16) \
2267  __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2268  } // namespace detail
2269 
2270 #define __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(base) \
2271  namespace detail { \
2272  __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 2) \
2273  __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 3) \
2274  __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 4) \
2275  __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 8) \
2276  __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 16) \
2277  __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2278  } // namespace detail
2279 
2280 #define __SYCL_DECLARE_FLOAT_VECTOR_CONVERTERS(base) \
2281  namespace detail { \
2282  __SYCL_DECLARE_FLOAT_CONVERTER(base, 2) \
2283  __SYCL_DECLARE_FLOAT_CONVERTER(base, 3) \
2284  __SYCL_DECLARE_FLOAT_CONVERTER(base, 4) \
2285  __SYCL_DECLARE_FLOAT_CONVERTER(base, 8) \
2286  __SYCL_DECLARE_FLOAT_CONVERTER(base, 16) \
2287  __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2288  } // namespace detail
2289 
2290 #define __SYCL_DECLARE_HALF_VECTOR_CONVERTERS(base) \
2291  namespace detail { \
2292  __SYCL_DECLARE_HALF_CONVERTER(base, 2) \
2293  __SYCL_DECLARE_HALF_CONVERTER(base, 3) \
2294  __SYCL_DECLARE_HALF_CONVERTER(base, 4) \
2295  __SYCL_DECLARE_HALF_CONVERTER(base, 8) \
2296  __SYCL_DECLARE_HALF_CONVERTER(base, 16) \
2297  template <> class BaseCLTypeConverter<base, 1> { \
2298  public: \
2299  using DataType = __half_t; \
2300  }; \
2301  } // namespace detail
2302 
2303 #define __SYCL_DECLARE_VECTOR_LONGLONG_CONVERTERS(base) \
2304  namespace detail { \
2305  __SYCL_DECLARE_LONGLONG_CONVERTER(base, 2) \
2306  __SYCL_DECLARE_LONGLONG_CONVERTER(base, 3) \
2307  __SYCL_DECLARE_LONGLONG_CONVERTER(base, 4) \
2308  __SYCL_DECLARE_LONGLONG_CONVERTER(base, 8) \
2309  __SYCL_DECLARE_LONGLONG_CONVERTER(base, 16) \
2310  template <> class BaseCLTypeConverter<base##long, 1> { \
2311  public: \
2312  using DataType = base##long; \
2313  }; \
2314  } // namespace detail
2315 
2316 #define __SYCL_DECLARE_SCHAR_VECTOR_CONVERTERS \
2317  namespace detail { \
2318  __SYCL_DECLARE_SCHAR_CONVERTER(2) \
2319  __SYCL_DECLARE_SCHAR_CONVERTER(3) \
2320  __SYCL_DECLARE_SCHAR_CONVERTER(4) \
2321  __SYCL_DECLARE_SCHAR_CONVERTER(8) \
2322  __SYCL_DECLARE_SCHAR_CONVERTER(16) \
2323  __SYCL_DECLARE_SCALAR_SCHAR_CONVERTER \
2324  } // namespace detail
2325 
2326 #define __SYCL_DECLARE_BOOL_VECTOR_CONVERTERS \
2327  namespace detail { \
2328  __SYCL_DECLARE_BOOL_CONVERTER(2) \
2329  __SYCL_DECLARE_BOOL_CONVERTER(3) \
2330  __SYCL_DECLARE_BOOL_CONVERTER(4) \
2331  __SYCL_DECLARE_BOOL_CONVERTER(8) \
2332  __SYCL_DECLARE_BOOL_CONVERTER(16) \
2333  __SYCL_DECLARE_SCALAR_BOOL_CONVERTER \
2334  } // namespace detail
2335 
2336 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2337 #define __SYCL_DECLARE_BYTE_VECTOR_CONVERTER \
2338  namespace detail { \
2339  __SYCL_DECLARE_BYTE_CONVERTER(2) \
2340  __SYCL_DECLARE_BYTE_CONVERTER(3) \
2341  __SYCL_DECLARE_BYTE_CONVERTER(4) \
2342  __SYCL_DECLARE_BYTE_CONVERTER(8) \
2343  __SYCL_DECLARE_BYTE_CONVERTER(16) \
2344  __SYCL_DECLARE_SCALAR_BYTE_CONVERTER \
2345  }
2346 #endif
2347 __SYCL_DECLARE_VECTOR_CONVERTERS(char)
2348 __SYCL_DECLARE_SCHAR_VECTOR_CONVERTERS
2349 __SYCL_DECLARE_BOOL_VECTOR_CONVERTERS
2350 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2351 __SYCL_DECLARE_BYTE_VECTOR_CONVERTER
2352 #endif
2353 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(uchar)
2354 __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(short)
2355 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(ushort)
2356 __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(int)
2357 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(uint)
2358 __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(long)
2359 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(ulong)
2360 __SYCL_DECLARE_VECTOR_LONGLONG_CONVERTERS(long)
2361 __SYCL_DECLARE_VECTOR_LONGLONG_CONVERTERS(ulong)
2362 __SYCL_DECLARE_HALF_VECTOR_CONVERTERS(half)
2363 __SYCL_DECLARE_FLOAT_VECTOR_CONVERTERS(float)
2364 __SYCL_DECLARE_FLOAT_VECTOR_CONVERTERS(double)
2365 
2366 #undef __SYCL_GET_CL_TYPE
2367 #undef __SYCL_GET_SCALAR_CL_TYPE
2368 #undef __SYCL_DECLARE_CONVERTER
2369 #undef __SYCL_DECLARE_VECTOR_CONVERTERS
2370 #undef __SYCL_DECLARE_SYCL_VEC
2371 #undef __SYCL_DECLARE_SYCL_VEC_WO_CONVERTERS
2372 #undef __SYCL_DECLARE_SCHAR_VECTOR_CONVERTERS
2373 #undef __SYCL_DECLARE_SCHAR_CONVERTER
2374 #undef __SYCL_DECLARE_SCALAR_SCHAR_CONVERTER
2375 #undef __SYCL_DECLARE_BOOL_VECTOR_CONVERTERS
2376 #undef __SYCL_DECLARE_BOOL_CONVERTER
2377 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2378 #undef __SYCL_DECLARE_BYTE_VECTOR_CONVERTER
2379 #undef __SYCL_DECLARE_BYTE_CONVERTER
2380 #undef __SYCL_DECLARE_SCALAR_BYTE_CONVERTER
2381 #endif
2382 #undef __SYCL_DECLARE_SCALAR_BOOL_CONVERTER
2383 #undef __SYCL_USE_EXT_VECTOR_TYPE__
2384 
2388 #define SYCL_DEVICE_COPYABLE 1
2389 
2397 template <typename T, typename = void>
2398 struct is_device_copyable : std::false_type {};
2399 
2400 template <typename T>
2401 struct is_device_copyable<
2402  T, std::enable_if_t<std::is_trivially_copyable<T>::value>>
2403  : std::true_type {};
2404 
2405 #if __cplusplus >= 201703L
2406 template <typename T>
2407 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
2408 #endif // __cplusplus >= 201703L
2409 
2410 // std::tuple<> is implicitly device copyable type.
2411 template <> struct is_device_copyable<std::tuple<>> : std::true_type {};
2412 
2413 // std::tuple<Ts...> is implicitly device copyable type if each type T of Ts...
2414 // is device copyable.
2415 template <typename T, typename... Ts>
2416 struct is_device_copyable<std::tuple<T, Ts...>>
2417  : detail::bool_constant<is_device_copyable<T>::value &&
2418  is_device_copyable<std::tuple<Ts...>>::value> {};
2419 
2420 // marray is device copyable if element type is device copyable and it is also
2421 // not trivially copyable (if the element type is trivially copyable, the marray
2422 // is device copyable by default).
2423 template <typename T, std::size_t N>
2424 struct is_device_copyable<
2425  sycl::marray<T, N>, std::enable_if_t<is_device_copyable<T>::value &&
2426  !std::is_trivially_copyable<T>::value>>
2427  : std::true_type {};
2428 
2429 namespace detail {
2430 template <typename T, typename = void>
2431 struct IsDeprecatedDeviceCopyable : std::false_type {};
2432 
2433 // TODO: using C++ attribute [[deprecated]] or the macro __SYCL2020_DEPRECATED
2434 // does not produce expected warning message for the type 'T'.
2435 template <typename T>
2436 struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
2437  IsDeprecatedDeviceCopyable<
2438  T, std::enable_if_t<std::is_trivially_copy_constructible<T>::value &&
2439  std::is_trivially_destructible<T>::value &&
2440  !is_device_copyable<T>::value>> : std::true_type {};
2441 
2442 #ifdef __SYCL_DEVICE_ONLY__
2443 // Checks that the fields of the type T with indices 0 to (NumFieldsToCheck - 1)
2444 // are device copyable.
2445 template <typename T, unsigned NumFieldsToCheck>
2446 struct CheckFieldsAreDeviceCopyable
2447  : CheckFieldsAreDeviceCopyable<T, NumFieldsToCheck - 1> {
2448  using FieldT = decltype(__builtin_field_type(T, NumFieldsToCheck - 1));
2449  static_assert(is_device_copyable<FieldT>::value ||
2450  detail::IsDeprecatedDeviceCopyable<FieldT>::value,
2451  "The specified type is not device copyable");
2452 };
2453 
2454 template <typename T> struct CheckFieldsAreDeviceCopyable<T, 0> {};
2455 
2456 // Checks that the base classes of the type T with indices 0 to
2457 // (NumFieldsToCheck - 1) are device copyable.
2458 template <typename T, unsigned NumBasesToCheck>
2459 struct CheckBasesAreDeviceCopyable
2460  : CheckBasesAreDeviceCopyable<T, NumBasesToCheck - 1> {
2461  using BaseT = decltype(__builtin_base_type(T, NumBasesToCheck - 1));
2462  static_assert(is_device_copyable<BaseT>::value ||
2463  detail::IsDeprecatedDeviceCopyable<BaseT>::value,
2464  "The specified type is not device copyable");
2465 };
2466 
2467 template <typename T> struct CheckBasesAreDeviceCopyable<T, 0> {};
2468 
2469 // All the captures of a lambda or functor of type FuncT passed to a kernel
2470 // must be is_device_copyable, which extends to bases and fields of FuncT.
2471 // Fields are captures of lambda/functors and bases are possible base classes
2472 // of functors also allowed by SYCL.
2473 // The SYCL-2020 implementation must check each of the fields & bases of the
2474 // type FuncT, only one level deep, which is enough to see if they are all
2475 // device copyable by using the result of is_device_copyable returned for them.
2476 // At this moment though the check also allowes using types for which
2477 // (is_trivially_copy_constructible && is_trivially_destructible) returns true
2478 // and (is_device_copyable) returns false. That is the deprecated behavior and
2479 // is currently/temporarily supported only to not break older SYCL programs.
2480 template <typename FuncT>
2481 struct CheckDeviceCopyable
2482  : CheckFieldsAreDeviceCopyable<FuncT, __builtin_num_fields(FuncT)>,
2483  CheckBasesAreDeviceCopyable<FuncT, __builtin_num_bases(FuncT)> {};
2484 
2485 // Below are two specializations for CheckDeviceCopyable when a kernel lambda
2486 // is wrapped after range rounding optimization.
2487 template <typename TransformedArgType, int Dims, typename KernelType>
2488 struct CheckDeviceCopyable<
2489  RoundedRangeKernel<TransformedArgType, Dims, KernelType>>
2490  : CheckDeviceCopyable<KernelType> {};
2491 
2492 template <typename TransformedArgType, int Dims, typename KernelType>
2493 struct CheckDeviceCopyable<
2494  RoundedRangeKernelWithKH<TransformedArgType, Dims, KernelType>>
2495  : CheckDeviceCopyable<KernelType> {};
2496 
2497 #endif // __SYCL_DEVICE_ONLY__
2498 } // namespace detail
2499 
2500 } // namespace sycl
2501 } // __SYCL_INLINE_NAMESPACE(cl)
2502 
2503 #undef __SYCL_ALIGNED_VAR
cl::sycl::detail::LessThan
Definition: types.hpp:185
cl::sycl::cl_long
std::int64_t cl_long
Definition: aliases.hpp:84
T
cl::sycl::detail::GetScalarOp::GetScalarOp
GetScalarOp(DataT Data)
Definition: types.hpp:148
cl::sycl::detail::RShift::operator()
constexpr T operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:204
type_traits.hpp
cl::sycl::ulong
unsigned long ulong
Definition: aliases.hpp:73
cl::sycl::trunc
detail::enable_if_t< detail::is_genfloat< T >::value, T > trunc(T x) __NOEXC
Definition: builtins.hpp:499
cl::sycl::detail::is_int_to_int
std::integral_constant< bool, std::is_integral< T >::value &&std::is_integral< R >::value > is_int_to_int
Definition: types.hpp:218
cl::sycl::detail::GreaterThan
Definition: types.hpp:179
cl::sycl::detail::is_uint_to_float
std::integral_constant< bool, std::is_unsigned< T >::value &&detail::is_floating_point< R >::value > is_uint_to_float
Definition: types.hpp:244
cl::sycl::rounding_mode::rtz
@ rtz
cl::sycl::detail::is_contained
Definition: type_list.hpp:54
cl::sycl::rounding_mode::rtn
@ rtn
cl::sycl::detail::LogicalAnd::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:192
cl::sycl::detail::RShift
Definition: types.hpp:203
cl::sycl::detail::LessEqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:174
cl::sycl::detail::vec_helper::RetType
T RetType
Definition: types.hpp:102
cl::sycl::detail::LShift::operator()
constexpr T operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:210
cl::sycl::floor
detail::enable_if_t< detail::is_genfloat< T >::value, T > floor(T x) __NOEXC
Definition: builtins.hpp:190
helpers.hpp
sycl
Definition: invoke_simd.hpp:68
cl::sycl::detail::is_floating_point
Definition: type_traits.hpp:223
cl::sycl::detail::rel_t
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
access.hpp
cl::sycl::detail::EqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:156
cl::sycl::detail::is_sint_to_float
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
cl::sycl::detail::GetOp::operator()
DataT operator()(DataT, DataT)
Definition: types.hpp:139
cl::sycl::detail::SwizzleOp
Definition: types.hpp:118
cl::sycl::vec_data_t
typename detail::vec_helper< T >::RetType vec_data_t
Definition: types.hpp:532
cl::sycl::detail::GreaterEqualTo
Definition: types.hpp:167
cl::sycl::detail::LogicalOr
Definition: types.hpp:197
cl::sycl::half
cl::sycl::detail::half_impl::half half
Definition: aliases.hpp:76
cl::sycl::detail::NotEqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:162
PI_ERROR_UNKNOWN
@ PI_ERROR_UNKNOWN
Definition: pi.h:131
cl::sycl::detail::GreaterThan::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:180
cl::sycl::detail::conditional_t
typename std::conditional< B, T, F >::type conditional_t
Definition: stl_type_traits.hpp:27
cl::sycl::detail::bool_constant
std::integral_constant< bool, V > bool_constant
Definition: stl_type_traits.hpp:40
generic_type_traits.hpp
cl::sycl::detail::BaseCLTypeConverter
Definition: types.hpp:120
cl::sycl::byte
unsigned char byte
Definition: image.hpp:59
cl::sycl::rounding_mode::rtp
@ rtp
multi_ptr.hpp
cl::sycl::detail::LShift
Definition: types.hpp:209
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::is_standard_type
std::integral_constant< bool, detail::is_sgentype< T >::value > is_standard_type
Definition: types.hpp:262
cl::sycl::detail::is_sint_to_from_uint
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
cl::sycl::cl_char
std::int8_t cl_char
Definition: aliases.hpp:78
cl::sycl::detail::GetOp
Definition: types.hpp:135
cl::sycl::detail::vec_helper
Definition: types.hpp:101
cl::sycl::image_channel_order::r
@ r
cl::sycl::detail::GetScalarOp::DataT
T DataT
Definition: types.hpp:147
cl::sycl::rounding_mode
rounding_mode
Definition: types.hpp:71
cl::sycl::detail::vec_helper::get
static constexpr RetType get(T value)
Definition: types.hpp:103
cl::sycl::detail::is_uint_to_uint
std::integral_constant< bool, is_sugeninteger< T >::value &&is_sugeninteger< R >::value > is_uint_to_uint
Definition: types.hpp:228
cl::sycl::detail::EqualTo
Definition: types.hpp:155
cl::sycl::elem
Definition: types.hpp:72
cl::sycl::image_channel_order::a
@ a
cl::sycl::detail::GetScalarOp::getValue
DataT getValue(size_t) const
Definition: types.hpp:149
cl::sycl::detail::GetOp::DataT
T DataT
Definition: types.hpp:137
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(cl::sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
cl::sycl::detail::LogicalOr::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:198
cl::sycl::ceil
detail::enable_if_t< detail::is_genfloat< T >::value, T > ceil(T x) __NOEXC
Definition: builtins.hpp:111
cl::sycl::uchar
unsigned char uchar
Definition: aliases.hpp:70
cl::sycl::detail::is_int_to_float
std::integral_constant< bool, std::is_integral< T >::value &&detail::is_floating_point< R >::value > is_int_to_float
Definition: types.hpp:249
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:82
cl::sycl::detail::is_sint_to_sint
std::integral_constant< bool, is_sigeninteger< T >::value &&is_sigeninteger< R >::value > is_sint_to_sint
Definition: types.hpp:223
cl::sycl::rint
detail::enable_if_t< detail::is_genfloat< T >::value, T > rint(T x) __NOEXC
Definition: builtins.hpp:415
std
Definition: accessor.hpp:2616
cl::sycl::rounding_mode::rte
@ rte
cl::sycl::detail::NotEqualTo
Definition: types.hpp:161
cl::sycl::detail::GetScalarOp
Definition: types.hpp:145
cl::sycl::detail::GetOp::getValue
DataT getValue(size_t) const
Definition: types.hpp:138
half_type.hpp
aliases.hpp
marray.hpp
cl::sycl::detail::convertImpl
detail::enable_if_t< is_float_to_int< T, R >::value, R > convertImpl(T Value)
Definition: types.hpp:288
common.hpp
cl::sycl::cl_short
std::int16_t cl_short
Definition: aliases.hpp:80
cl::sycl::detail::LessEqualTo
Definition: types.hpp:173
cl::sycl::ushort
unsigned short ushort
Definition: aliases.hpp:71
cl::sycl::detail::LogicalAnd
Definition: types.hpp:191
cl::sycl::detail::is_float_to_float
std::integral_constant< bool, detail::is_floating_point< T >::value &&detail::is_floating_point< R >::value > is_float_to_float
Definition: types.hpp:259
cl::sycl::uint
unsigned int uint
Definition: aliases.hpp:72
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::detail::LessThan::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:186
cl::sycl::detail::GreaterEqualTo::operator()
constexpr rel_t< T > operator()(const T &Lhs, const T &Rhs) const
Definition: types.hpp:168
cl::sycl::detail::is_float_to_int
std::integral_constant< bool, detail::is_floating_point< T >::value &&std::is_integral< R >::value > is_float_to_int
Definition: types.hpp:254
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12