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