15 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
17 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
22 #ifndef __has_extension
23 #define __has_extension(x) 0
25 #ifdef __HAS_EXT_VECTOR_TYPE__
26 #error "Undefine __HAS_EXT_VECTOR_TYPE__ macro"
28 #if __has_extension(attribute_ext_vector_type)
29 #define __HAS_EXT_VECTOR_TYPE__
33 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
35 #if !defined(__HAS_EXT_VECTOR_TYPE__) && defined(__SYCL_DEVICE_ONLY__)
36 #error "SYCL device compiler is built without ext_vector_type support"
39 #if defined(__SYCL_DEVICE_ONLY__)
40 #define __SYCL_USE_EXT_VECTOR_TYPE__
43 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
45 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
47 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
48 #error "Undefine __SYCL_USE_EXT_VECTOR_TYPE__ macro"
50 #ifdef __HAS_EXT_VECTOR_TYPE__
51 #if defined(__SYCL_DEVICE_ONLY__) || !defined(__NO_EXT_VECTOR_TYPE_ON_HOST__)
52 #define __SYCL_USE_EXT_VECTOR_TYPE__
54 #elif defined(__SYCL_DEVICE_ONLY__)
57 #error "SYCL device compiler is built without ext_vector_type support"
58 #endif // __HAS_EXT_VECTOR_TYPE__
60 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
87 #include <type_traits>
91 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
93 #ifndef __SYCL_USE_EXT_VECTOR_TYPE__
98 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
104 inline namespace _V1 {
107 static constexpr
int x = 0;
108 static constexpr
int y = 1;
109 static constexpr
int z = 2;
110 static constexpr
int w = 3;
111 static constexpr
int r = 0;
112 static constexpr
int g = 1;
113 static constexpr
int b = 2;
114 static constexpr
int a = 3;
115 static constexpr
int s0 = 0;
116 static constexpr
int s1 = 1;
117 static constexpr
int s2 = 2;
118 static constexpr
int s3 = 3;
119 static constexpr
int s4 = 4;
120 static constexpr
int s5 = 5;
121 static constexpr
int s6 = 6;
122 static constexpr
int s7 = 7;
123 static constexpr
int s8 = 8;
124 static constexpr
int s9 = 9;
125 static constexpr
int sA = 10;
126 static constexpr
int sB = 11;
127 static constexpr
int sC = 12;
128 static constexpr
int sD = 13;
129 static constexpr
int sE = 14;
130 static constexpr
int sF = 15;
136 template <
typename _IN,
typename T8,
typename T16,
typename T32,
typename T64>
138 sizeof(_IN) == 1, T8,
139 std::conditional_t<
sizeof(_IN) == 2, T16,
140 std::conditional_t<
sizeof(_IN) == 4, T32, T64>>>;
152 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
162 template <
typename VecT,
typename OperationLeftT,
typename OperationRightT,
163 template <
typename>
class OperationCurrentT,
int... Indexes>
166 template <
typename T,
int N,
typename V =
void>
struct VecStorage;
169 template <
typename DataT>
170 using rel_t =
typename std::conditional_t<
172 typename std::conditional_t<
174 typename std::conditional_t<
176 typename std::conditional_t<
sizeof(DataT) ==
204 return (Lhs == Rhs) ? -1 : 0;
210 return (Lhs != Rhs) ? -1 : 0;
216 return (Lhs >= Rhs) ? -1 : 0;
222 return (Lhs <= Rhs) ? -1 : 0;
228 return (Lhs > Rhs) ? -1 : 0;
234 return (Lhs < Rhs) ? -1 : 0;
240 return (Lhs && Rhs) ? -1 : 0;
246 return (Lhs || Rhs) ? -1 : 0;
263 template <
typename TransformedArgType,
int Dims,
typename KernelType>
264 class RoundedRangeKernel;
265 template <
typename TransformedArgType,
int Dims,
typename KernelType>
266 class RoundedRangeKernelWithKH;
272 template <
typename T>
275 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
277 #if defined(_WIN32) && (_MSC_VER)
288 #pragma message("Alignment of class vec is not in accordance with SYCL \
289 specification requirements, a limitation of the MSVC compiler(Error C2719).\
290 Requested alignment applied, limited at 64.")
291 #define __SYCL_ALIGNED_VAR(type, x, var) \
292 type __declspec(align((x < 64) ? x : 64)) var
294 #define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
299 template <
typename Type,
int NumElements>
class vec {
308 using DataType =
typename detail::VecStorage<DataT, NumElements>::DataType;
310 static constexpr
bool IsHostHalf =
311 std::is_same_v<DataT, sycl::detail::half_impl::half> &&
315 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
317 static constexpr
size_t AdjustedNum = (NumElements == 3) ? 4 : NumElements;
318 static constexpr
size_t Sz =
sizeof(DataT) * AdjustedNum;
319 static constexpr
bool IsSizeGreaterThanMaxAlign =
320 (Sz > detail::MaxVecAlignment);
327 static constexpr
bool IsUsingArrayOnDevice =
328 (IsHostHalf || IsSizeGreaterThanMaxAlign);
330 #if defined(__SYCL_DEVICE_ONLY__)
331 static constexpr
bool NativeVec = NumElements > 1 && !IsUsingArrayOnDevice;
332 static constexpr
bool IsUsingArrayOnHost =
false;
334 static constexpr
bool NativeVec =
false;
335 static constexpr
bool IsUsingArrayOnHost =
true;
338 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
340 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
341 static constexpr
bool IsUsingArrayOnDevice = IsHostHalf;
342 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
344 static constexpr
int getNumElements() {
return NumElements; }
347 template <
int Counter,
int MaxValue,
class...>
348 struct SizeChecker : std::conditional_t<Counter == MaxValue, std::true_type,
351 template <
int Counter,
int MaxValue,
typename DataT_,
class... tail>
352 struct SizeChecker<Counter, MaxValue, DataT_, tail...>
353 : std::conditional_t<Counter + 1 <= MaxValue,
354 SizeChecker<Counter + 1, MaxValue, tail...>,
358 template <typename DataT_, typename T, std::size_t... Is>
359 static constexpr std::array<DataT_, sizeof...(Is)>
360 VecToArray(const vec<T, sizeof...(Is)> &V, std::index_sequence<Is...>) {
361 return {static_cast<DataT_>(V.getValue(Is))...};
363 template <typename DataT_, typename T, int N, typename T2, typename T3,
364 template <typename> class T4, int... T5, std::size_t... Is>
365 static constexpr std::array<DataT_, sizeof...(Is)>
366 VecToArray(const detail::SwizzleOp<vec<T, N>, T2, T3, T4, T5...> &V,
367 std::index_sequence<Is...>) {
368 return {static_cast<DataT_>(V.getValue(Is))...};
370 template <typename DataT_, typename T, int N, typename T2, typename T3,
371 template <typename> class T4, int... T5, std::size_t... Is>
372 static constexpr std::array<DataT_, sizeof...(Is)>
373 VecToArray(const detail::SwizzleOp<const vec<T, N>, T2, T3, T4, T5...> &V,
374 std::index_sequence<Is...>) {
375 return {static_cast<DataT_>(V.getValue(Is))...};
377 template <typename DataT_, typename T, int N>
378 static constexpr std::array<DataT_, N>
379 FlattenVecArgHelper(const vec<T, N> &A) {
380 return VecToArray<DataT_>(A, std::make_index_sequence<N>());
382 template <typename DataT_, typename T, int N, typename T2, typename T3,
383 template <typename> class T4, int... T5>
384 static constexpr std::array<DataT_, sizeof...(T5)> FlattenVecArgHelper(
385 const detail::SwizzleOp<vec<T, N>, T2, T3, T4, T5...> &A) {
386 return VecToArray<DataT_>(A, std::make_index_sequence<sizeof...(T5)>());
388 template <typename DataT_, typename T, int N, typename T2, typename T3,
389 template <typename> class T4, int... T5>
390 static constexpr std::array<DataT_, sizeof...(T5)> FlattenVecArgHelper(
391 const detail::SwizzleOp<const vec<T, N>, T2, T3, T4, T5...> &A) {
392 return VecToArray<DataT_>(A, std::make_index_sequence<sizeof...(T5)>());
394 template <typename DataT_, typename T>
395 static constexpr auto FlattenVecArgHelper(const T &A) {
396 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
397 return std::array<DataT_, 1>{vec_data<DataT_>::get(static_cast<DataT_>(A))};
399 return std::array<DataT_, 1>{vec_data<DataT_>::get(A)};
402 template <typename DataT_, typename T> struct FlattenVecArg {
403 constexpr auto operator()(const T &A) const {
404 return FlattenVecArgHelper<DataT_>(A);
409 template <typename DataT_, typename... ArgTN>
410 using VecArgArrayCreator =
411 detail::ArrayCreator<DataT_, FlattenVecArg, ArgTN...>;
413 #define __SYCL_ALLOW_VECTOR_SIZES(num_elements) \
414 template <int Counter, int MaxValue, typename DataT_, class... tail> \
415 struct SizeChecker<Counter, MaxValue, vec<DataT_, num_elements>, tail...> \
416 : std::conditional_t< \
417 Counter + (num_elements) <= MaxValue, \
418 SizeChecker<Counter + (num_elements), MaxValue, tail...>, \
419 std::false_type> {}; \
420 template <int Counter, int MaxValue, typename DataT_, typename T2, \
421 typename T3, template <typename> class T4, int... T5, \
423 struct SizeChecker< \
425 detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
427 : std::conditional_t< \
428 Counter + sizeof...(T5) <= MaxValue, \
429 SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
430 std::false_type> {}; \
431 template <int Counter, int MaxValue, typename DataT_, typename T2, \
432 typename T3, template <typename> class T4, int... T5, \
434 struct SizeChecker< \
436 detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
438 : std::conditional_t< \
439 Counter + sizeof...(T5) <= MaxValue, \
440 SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
443 __SYCL_ALLOW_VECTOR_SIZES(1)
444 __SYCL_ALLOW_VECTOR_SIZES(2)
445 __SYCL_ALLOW_VECTOR_SIZES(3)
446 __SYCL_ALLOW_VECTOR_SIZES(4)
447 __SYCL_ALLOW_VECTOR_SIZES(8)
448 __SYCL_ALLOW_VECTOR_SIZES(16)
449 #undef __SYCL_ALLOW_VECTOR_SIZES
452 template <typename T, typename DataT_>
453 struct TypeChecker : std::is_convertible<T, DataT_> {};
454 #define __SYCL_ALLOW_VECTOR_TYPES(num_elements) \
455 template <typename DataT_> \
456 struct TypeChecker<vec<DataT_, num_elements>, DataT_> : std::true_type {}; \
457 template <typename DataT_, typename T2, typename T3, \
458 template <typename> class T4, int... T5> \
459 struct TypeChecker< \
460 detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, DataT_> \
461 : std::true_type {}; \
462 template <typename DataT_, typename T2, typename T3, \
463 template <typename> class T4, int... T5> \
464 struct TypeChecker< \
465 detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
466 DataT_> : std::true_type {};
468 __SYCL_ALLOW_VECTOR_TYPES(1)
469 __SYCL_ALLOW_VECTOR_TYPES(2)
470 __SYCL_ALLOW_VECTOR_TYPES(3)
471 __SYCL_ALLOW_VECTOR_TYPES(4)
472 __SYCL_ALLOW_VECTOR_TYPES(8)
473 __SYCL_ALLOW_VECTOR_TYPES(16)
474 #undef __SYCL_ALLOW_VECTOR_TYPES
476 template <int... Indexes>
478 detail::SwizzleOp<vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
479 detail::GetOp, Indexes...>;
481 template <int... Indexes>
483 detail::SwizzleOp<const vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
484 detail::GetOp, Indexes...>;
487 template <typename... argTN>
488 using EnableIfSuitableTypes = typename std::enable_if_t<
489 std::conjunction_v<TypeChecker<argTN, DataT>...>>;
491 template <typename... argTN>
492 using EnableIfSuitableNumElements =
493 typename std::enable_if_t<SizeChecker<0, NumElements, argTN...>::value>;
495 template <size_t... Is>
496 constexpr vec(const std::array<vec_data_t<DataT>, NumElements> &Arr,
497 std::index_sequence<Is...>)
498 : m_Data{vec_data_t<DataT>(static_cast<DataT>(Arr[Is]))...} {}
501 using element_type = DataT;
502 using rel_t = detail::rel_t<DataT>;
504 #ifdef __SYCL_DEVICE_ONLY__
505 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
507 typename detail::VecStorage<DataT, NumElements>::VectorDataType;
509 using vector_t = DataType;
515 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
516 constexpr vec(const vec &Rhs) = default;
517 constexpr vec(vec &&Rhs) = default;
519 constexpr vec &operator=(const vec &Rhs) = default;
522 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
525 #ifdef __SYCL_DEVICE_ONLY__
526 vec(const vec &Rhs) = default;
527 vec &operator=(const vec &Rhs) = default;
529 constexpr vec(const vec &Rhs) : m_Data(Rhs.m_Data) {}
530 constexpr vec &operator=(const vec &Rhs) = default;
533 vec(vec &&Rhs) = default;
537 template <typename Ty = DataT>
538 typename std::enable_if_t<!std::is_same_v<Ty, rel_t> &&
539 std::is_convertible_v<vec_data_t<Ty>, rel_t>,
541 operator=(const vec<rel_t, NumElements> &Rhs) {
542 *this = Rhs.template as<vec>();
546 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
547 template <typename T = void>
548 using EnableIfNotHostHalf = typename std::enable_if_t<!IsHostHalf, T>;
550 template <typename T = void>
551 using EnableIfHostHalf = typename std::enable_if_t<IsHostHalf, T>;
553 template <typename T = void>
554 using EnableIfUsingArrayOnDevice =
555 typename std::enable_if_t<IsUsingArrayOnDevice, T>;
557 template <typename T = void>
558 using EnableIfNotUsingArrayOnDevice =
559 typename std::enable_if_t<!IsUsingArrayOnDevice, T>;
562 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
564 template <typename T = void>
565 using EnableIfUsingArray =
566 typename std::enable_if_t<IsUsingArrayOnDevice || IsUsingArrayOnHost, T>;
568 template <typename T = void>
569 using EnableIfNotUsingArray =
570 typename std::enable_if_t<!IsUsingArrayOnDevice && !IsUsingArrayOnHost,
574 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
576 template <typename Ty = DataT>
577 explicit constexpr vec(const EnableIfNotUsingArrayOnDevice<Ty> &arg)
578 : m_Data{DataType(vec_data<Ty>::get(arg))} {}
580 template <typename Ty = DataT>
581 typename std::enable_if_t<
582 std::is_fundamental_v<vec_data_t<Ty>> ||
583 std::is_same_v<typename std::remove_const_t<Ty>, half>,
585 operator=(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) {
586 m_Data = (DataType)vec_data<Ty>::get(Rhs);
590 template <typename Ty = DataT>
591 explicit constexpr vec(const EnableIfUsingArrayOnDevice<Ty> &arg)
592 : vec{detail::RepeatValue<NumElements>(
593 static_cast<vec_data_t<DataT>>(arg)),
594 std::make_index_sequence<NumElements>()} {}
596 template <typename Ty = DataT>
597 typename std::enable_if_t<
598 std::is_fundamental_v<vec_data_t<Ty>> ||
599 std::is_same_v<typename std::remove_const_t<Ty>, half>,
601 operator=(const EnableIfUsingArrayOnDevice<Ty> &Rhs) {
602 for (int i = 0; i < NumElements; ++i) {
608 explicit constexpr vec(const DataT &arg)
609 : vec{detail::RepeatValue<NumElements>(
610 static_cast<vec_data_t<DataT>>(arg)),
611 std::make_index_sequence<NumElements>()} {}
613 template <typename Ty = DataT>
614 typename std::enable_if_t<
615 std::is_fundamental_v<vec_data_t<Ty>> ||
616 std::is_same_v<typename std::remove_const_t<Ty>, half>,
618 operator=(const DataT &Rhs) {
619 for (int i = 0; i < NumElements; ++i) {
626 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
632 template <int IdxNum, typename T = void>
633 using EnableIfMultipleElems = typename std::enable_if_t<
634 std::is_convertible_v<T, DataT> && NumElements == IdxNum, DataT>;
635 template <typename Ty = DataT>
636 constexpr vec(const EnableIfMultipleElems<2, Ty> Arg0,
637 const EnableIfNotUsingArrayOnDevice<Ty> Arg1)
638 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1)} {}
639 template <typename Ty = DataT>
640 constexpr vec(const EnableIfMultipleElems<3, Ty> Arg0,
641 const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2)
642 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
643 vec_data<Ty>::get(Arg2)} {}
644 template <typename Ty = DataT>
645 constexpr vec(const EnableIfMultipleElems<4, Ty> Arg0,
646 const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2,
648 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
649 vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3)} {}
650 template <typename Ty = DataT>
651 constexpr vec(const EnableIfMultipleElems<8, Ty> Arg0,
652 const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2,
653 const DataT Arg3, const DataT Arg4, const DataT Arg5,
654 const DataT Arg6, const DataT Arg7)
655 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
656 vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3),
657 vec_data<Ty>::get(Arg4), vec_data<Ty>::get(Arg5),
658 vec_data<Ty>::get(Arg6), vec_data<Ty>::get(Arg7)} {}
659 template <typename Ty = DataT>
660 constexpr vec(const EnableIfMultipleElems<16, Ty> Arg0,
661 const EnableIfNotUsingArrayOnDevice<Ty> Arg1, const DataT Arg2,
662 const DataT Arg3, const DataT Arg4, const DataT Arg5,
663 const DataT Arg6, const DataT Arg7, const DataT Arg8,
664 const DataT Arg9, const DataT ArgA, const DataT ArgB,
665 const DataT ArgC, const DataT ArgD, const DataT ArgE,
667 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
668 vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3),
669 vec_data<Ty>::get(Arg4), vec_data<Ty>::get(Arg5),
670 vec_data<Ty>::get(Arg6), vec_data<Ty>::get(Arg7),
671 vec_data<Ty>::get(Arg8), vec_data<Ty>::get(Arg9),
672 vec_data<Ty>::get(ArgA), vec_data<Ty>::get(ArgB),
673 vec_data<Ty>::get(ArgC), vec_data<Ty>::get(ArgD),
674 vec_data<Ty>::get(ArgE), vec_data<Ty>::get(ArgF)} {}
679 template <typename... argTN, typename = EnableIfSuitableTypes<argTN...>,
680 typename = EnableIfSuitableNumElements<argTN...>>
681 constexpr vec(const argTN &...args)
682 : vec{VecArgArrayCreator<vec_data_t<DataT>, argTN...>::Create(args...),
683 std::make_index_sequence<NumElements>()} {}
685 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
687 #ifdef __SYCL_DEVICE_ONLY__
688 template <typename vector_t_ = vector_t,
690 typename std::enable_if_t<std::is_same_v<vector_t_, vector_t> &&
691 !std::is_same_v<vector_t_, DataT>>>
692 constexpr vec(vector_t openclVector) {
693 if constexpr (!IsUsingArrayOnDevice) {
694 m_Data = openclVector;
696 m_Data = bit_cast<DataType>(openclVector);
700 operator vector_t() const {
701 if constexpr (!IsUsingArrayOnDevice) {
704 auto ptr = bit_cast<const vector_t *>((&m_Data)->data());
712 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
714 #ifdef __SYCL_DEVICE_ONLY__
715 template <typename vector_t_ = vector_t,
717 typename std::enable_if_t<std::is_same_v<vector_t_, vector_t> &&
718 !std::is_same_v<vector_t_, DataT>>>
719 constexpr vec(vector_t openclVector) : m_Data(openclVector) {}
720 operator vector_t() const { return m_Data; }
726 template <int N = NumElements>
727 operator typename std::enable_if_t<N == 1, DataT>() const {
728 return vec_data<DataT>::get(m_Data);
731 __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
732 static constexpr size_t get_count() { return size(); }
733 static constexpr size_t size() noexcept { return NumElements; }
734 __SYCL2020_DEPRECATED(
735 "get_size() is deprecated, please use byte_size() instead")
736 static constexpr size_t get_size() { return byte_size(); }
737 static constexpr size_t byte_size() noexcept { return sizeof(m_Data); }
741 template <typename convertT,
742 rounding_mode roundingMode = rounding_mode::automatic>
744 std::is_same_v<vec_data_t<DataT>, vec_data_t<convertT>> ||
745 std::is_same_v<detail::ConvertToOpenCLType_t<vec_data_t<DataT>>,
746 detail::ConvertToOpenCLType_t<vec_data_t<convertT>>>,
747 vec<convertT, NumElements>>
749 static_assert(std::is_integral_v<vec_data_t<convertT>> ||
750 detail::is_floating_point<convertT>::value,
751 "Unsupported convertT");
752 if constexpr (!std::is_same_v<DataT, convertT>) {
754 vec<convertT, NumElements> Result;
755 for (size_t I = 0; I < NumElements; ++I)
756 Result.setValue(I, static_cast<convertT>(getValue(I)));
765 template <typename convertT,
766 rounding_mode roundingMode = rounding_mode::automatic>
768 !std::is_same_v<vec_data_t<DataT>, vec_data_t<convertT>> &&
769 !std::is_same_v<detail::ConvertToOpenCLType_t<vec_data_t<DataT>>,
770 detail::ConvertToOpenCLType_t<vec_data_t<convertT>>>,
771 vec<convertT, NumElements>>
773 static_assert(std::is_integral_v<vec_data_t<convertT>> ||
774 detail::is_floating_point<convertT>::value,
775 "Unsupported convertT");
776 using T = vec_data_t<DataT>;
777 using R = vec_data_t<convertT>;
778 using OpenCLT = detail::ConvertToOpenCLType_t<T>;
779 using OpenCLR = detail::ConvertToOpenCLType_t<R>;
780 vec<convertT, NumElements> Result;
782 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
783 using OpenCLVecT = OpenCLT __attribute__((ext_vector_type(NumElements)));
784 using OpenCLVecR = OpenCLR __attribute__((ext_vector_type(NumElements)));
786 constexpr bool canUseNativeVectorConvert =
792 NativeVec && vec<convertT, NumElements>::NativeVec &&
796 std::is_convertible_v<decltype(m_Data), OpenCLVecT> &&
797 std::is_convertible_v<decltype(Result.m_Data), OpenCLVecR> &&
800 !detail::is_sint_to_from_uint<T, R>::value &&
806 !std::is_same_v<convertT, bool>;
807 if constexpr (canUseNativeVectorConvert) {
808 Result.m_Data = detail::convertImpl<T, R, roundingMode, NumElements,
809 OpenCLVecT, OpenCLVecR>(m_Data);
815 for (size_t I = 0; I < NumElements; ++I) {
817 I, vec_data<convertT>::get(
818 detail::convertImpl<T, R, roundingMode, 1, OpenCLT, OpenCLR>(
819 vec_data<DataT>::get(getValue(I)))));
826 template <typename asT> asT as() const {
827 static_assert((sizeof(*this) == sizeof(asT)),
828 "The new SYCL vec type must have the same storage size in "
829 "bytes as this SYCL vec");
831 detail::is_contained<asT, detail::gtl::vector_basic_list>::value ||
832 detail::is_contained<asT, detail::gtl::vector_bool_list>::value,
833 "asT must be SYCL vec of a different element type and "
834 "number of elements specified by asT");
836 detail::memcpy(&Result.m_Data, &m_Data, sizeof(decltype(Result.m_Data)));
840 template <int... SwizzleIndexes> Swizzle<SwizzleIndexes...> swizzle() {
844 template <int... SwizzleIndexes>
845 ConstSwizzle<SwizzleIndexes...> swizzle() const {
861 const DataT &operator[](int i) const {
862 return reinterpret_cast<const DataT *>(&m_Data)[i];
865 DataT &operator[](int i) { return reinterpret_cast<DataT *>(&m_Data)[i]; }
874 template <int Index> struct Indexer {
875 static constexpr int value = Index;
879 #ifdef __SYCL_ACCESS_RETURN
880 #error "Undefine __SYCL_ACCESS_RETURN macro"
882 #define __SYCL_ACCESS_RETURN this
883 #include "swizzles.def"
884 #undef __SYCL_ACCESS_RETURN
887 template <access::address_space Space, access::decorated DecorateAddress>
888 void load(size_t Offset, multi_ptr<const DataT, Space, DecorateAddress> Ptr) {
889 for (int I = 0; I < NumElements; I++) {
890 setValue(I, *multi_ptr<const DataT, Space, DecorateAddress>(
891 Ptr + Offset * NumElements + I));
894 template <access::address_space Space, access::decorated DecorateAddress>
895 void load(size_t Offset, multi_ptr<DataT, Space, DecorateAddress> Ptr) {
896 multi_ptr<const DataT, Space, DecorateAddress> ConstPtr(Ptr);
897 load(Offset, ConstPtr);
899 template <int Dimensions, access::mode Mode,
900 access::placeholder IsPlaceholder, access::target Target,
901 typename PropertyListT>
904 accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
906 multi_ptr<const DataT, detail::TargetToAS<Target>::AS,
907 access::decorated::yes>
909 load(Offset, MultiPtr);
911 template <access::address_space Space, access::decorated DecorateAddress>
912 void store(size_t Offset,
913 multi_ptr<DataT, Space, DecorateAddress> Ptr) const {
914 for (int I = 0; I < NumElements; I++) {
915 *multi_ptr<DataT, Space, DecorateAddress>(Ptr + Offset * NumElements +
919 template <int Dimensions, access::mode Mode,
920 access::placeholder IsPlaceholder, access::target Target,
921 typename PropertyListT>
924 accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
926 multi_ptr<DataT, detail::TargetToAS<Target>::AS, access::decorated::yes>
928 store(Offset, MultiPtr);
931 void ConvertToDataT() {
932 for (size_t i = 0; i < NumElements; ++i) {
933 DataT tmp = getValue(i);
939 #error "Undefine __SYCL_BINOP macro"
942 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
943 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \
944 template <typename Ty = vec> \
945 vec operator BINOP(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) const { \
947 Ret.m_Data = m_Data BINOP Rhs.m_Data; \
948 if constexpr (std::is_same_v<Type, bool> && CONVERT) { \
949 Ret.ConvertToDataT(); \
953 template <typename Ty = vec> \
954 vec operator BINOP(const EnableIfUsingArrayOnDevice<Ty> &Rhs) const { \
956 for (size_t I = 0; I < NumElements; ++I) { \
957 Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
961 template <typename T> \
962 typename std::enable_if_t< \
963 std::is_convertible_v<DataT, T> && \
964 (std::is_fundamental_v<vec_data_t<T>> || \
965 std::is_same_v<typename std::remove_const_t<T>, half>), \
967 operator BINOP(const T & Rhs) const { \
968 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
970 vec &operator OPASSIGN(const vec & Rhs) { \
971 *this = *this BINOP Rhs; \
974 template <int Num = NumElements> \
975 typename std::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
976 const DataT & Rhs) { \
977 *this = *this BINOP vec(Rhs); \
982 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
983 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \
984 vec operator BINOP(const vec &Rhs) const { \
986 if constexpr (NativeVec) \
987 Ret.m_Data = m_Data BINOP Rhs.m_Data; \
989 for (size_t I = 0; I < NumElements; ++I) \
990 Ret.setValue(I, (DataT)(vec_data<DataT>::get(getValue( \
991 I)) BINOP vec_data<DataT>::get(Rhs.getValue(I)))); \
994 template <typename T> \
995 typename std::enable_if_t< \
996 std::is_convertible_v<DataT, T> && \
997 (std::is_fundamental_v<vec_data_t<T>> || \
998 std::is_same_v<typename std::remove_const_t<T>, half>), \
1000 operator BINOP(const T & Rhs) const { \
1001 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1003 vec &operator OPASSIGN(const vec & Rhs) { \
1004 *this = *this BINOP Rhs; \
1007 template <int Num = NumElements> \
1008 typename std::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1009 const DataT & Rhs) { \
1010 *this = *this BINOP vec(Rhs); \
1015 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1016 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \
1017 vec operator BINOP(const vec &Rhs) const { \
1019 for (size_t I = 0; I < NumElements; ++I) { \
1020 Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
1024 template <typename T> \
1025 typename std::enable_if_t< \
1026 std::is_convertible_v<DataT, T> && \
1027 (std::is_fundamental_v<vec_data_t<T>> || \
1028 std::is_same_v<typename std::remove_const_t<T>, half>), \
1030 operator BINOP(const T & Rhs) const { \
1031 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1033 vec &operator OPASSIGN(const vec & Rhs) { \
1034 *this = *this BINOP Rhs; \
1037 template <int Num = NumElements> \
1038 typename std::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1039 const DataT & Rhs) { \
1040 *this = *this BINOP vec(Rhs); \
1047 __SYCL_BINOP(+, +=, true)
1048 __SYCL_BINOP(-, -=, true)
1049 __SYCL_BINOP(*, *=, false)
1050 __SYCL_BINOP(/, /=, false)
1054 __SYCL_BINOP(%, %=, false)
1055 __SYCL_BINOP(|, |=, false)
1056 __SYCL_BINOP(&, &=, false)
1057 __SYCL_BINOP(^, ^=, false)
1058 __SYCL_BINOP(>>, >>=, false)
1059 __SYCL_BINOP(<<, <<=, true)
1061 #undef __SYCL_BINOP_HELP
1070 #ifdef __SYCL_RELLOGOP
1071 #error "Undefine __SYCL_RELLOGOP macro"
1075 #ifdef __SYCL_DEVICE_ONLY__
1076 #define __SYCL_RELLOGOP(RELLOGOP) \
1077 vec<rel_t, NumElements> operator RELLOGOP(const vec & Rhs) const { \
1078 vec<rel_t, NumElements> Ret{}; \
1081 if constexpr (IsUsingArrayOnDevice && \
1082 (std::string_view(#RELLOGOP) == "||" || \
1083 std::string_view(#RELLOGOP) == "&&")) { \
1084 for (size_t I = 0; I < NumElements; ++I) { \
1086 -(vec_data<DataT>::get(getValue(I)) \
1087 RELLOGOP vec_data<DataT>::get(Rhs.getValue(I)))); \
1090 Ret = vec<rel_t, NumElements>( \
1091 (typename vec<rel_t, NumElements>::vector_t)( \
1092 m_Data RELLOGOP Rhs.m_Data)); \
1093 if (NumElements == 1) \
1098 template <typename T> \
1099 typename std::enable_if_t<std::is_convertible_v<T, DataT> && \
1100 (std::is_fundamental_v<vec_data_t<T>> || \
1101 std::is_same_v<T, half>), \
1102 vec<rel_t, NumElements>> \
1103 operator RELLOGOP(const T & Rhs) const { \
1104 return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1107 #define __SYCL_RELLOGOP(RELLOGOP) \
1108 vec<rel_t, NumElements> operator RELLOGOP(const vec & Rhs) const { \
1109 vec<rel_t, NumElements> Ret{}; \
1110 for (size_t I = 0; I < NumElements; ++I) { \
1111 Ret.setValue(I, -(vec_data<DataT>::get(getValue(I)) \
1112 RELLOGOP vec_data<DataT>::get(Rhs.getValue(I)))); \
1116 template <typename T> \
1117 typename std::enable_if_t<std::is_convertible_v<T, DataT> && \
1118 (std::is_fundamental_v<vec_data_t<T>> || \
1119 std::is_same_v<T, half>), \
1120 vec<rel_t, NumElements>> \
1121 operator RELLOGOP(const T & Rhs) const { \
1122 return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1135 #undef __SYCL_RELLOGOP
1138 #error "Undefine __SYCL_UOP macro"
1140 #define __SYCL_UOP(UOP, OPASSIGN) \
1141 vec &operator UOP() { \
1142 *this OPASSIGN vec_data<DataT>::get(1); \
1145 vec operator UOP(int) { \
1147 *this OPASSIGN vec_data<DataT>::get(1); \
1155 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1159 template <typename T = DataT>
1160 typename std::enable_if_t<!std::is_floating_point_v<vec_data_t<T>> &&
1161 (!IsUsingArrayOnDevice && !IsUsingArrayOnHost),
1164 vec Ret{(typename vec::DataType) ~m_Data};
1165 if constexpr (std::is_same_v<Type, bool>) {
1166 Ret.ConvertToDataT();
1170 template <typename T = DataT>
1171 typename std::enable_if_t<!std::is_floating_point_v<vec_data_t<T>> &&
1172 (IsUsingArrayOnDevice || IsUsingArrayOnHost),
1176 for (size_t I = 0; I < NumElements; ++I) {
1177 Ret.setValue(I, ~getValue(I));
1182 template <typename T>
1183 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1184 using OpNotRet = detail::rel_t<T>;
1190 template <typename T = DataT, int N = NumElements>
1191 EnableIfNotUsingArray<vec<OpNotRet<T>, N>> operator!() const {
1192 return vec<T, N>{(typename vec<DataT, NumElements>::DataType) !m_Data}
1193 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1194 .template as<vec<OpNotRet<T>, N>>();
1202 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1203 template <typename T = DataT, int N = NumElements>
1204 typename std::enable_if_t<std::is_same_v<std::byte, T> &&
1205 (IsUsingArrayOnDevice || IsUsingArrayOnHost),
1206 vec<OpNotRet<T>, N>>
1209 for (size_t I = 0; I < NumElements; ++I) {
1210 Ret.setValue(I, std::byte{!vec_data<DataT>::get(getValue(I))});
1212 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1213 return Ret.template as<vec<OpNotRet<T>, N>>();
1219 template <typename T = DataT, int N = NumElements>
1220 typename std::enable_if_t<!std::is_same_v<std::byte, T> &&
1221 (IsUsingArrayOnDevice || IsUsingArrayOnHost),
1222 vec<OpNotRet<T>, N>>
1225 for (size_t I = 0; I < NumElements; ++I)
1226 Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1227 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1228 return Ret.template as<vec<OpNotRet<T>, N>>();
1234 template <typename T = DataT, int N = NumElements>
1235 EnableIfUsingArray<vec<OpNotRet<T>, N>> operator!() const {
1237 for (size_t I = 0; I < NumElements; ++I)
1238 Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1239 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1240 return Ret.template as<vec<OpNotRet<T>, N>>();
1248 template <typename T = vec> EnableIfNotUsingArray<T> operator+() const {
1249 return vec{+m_Data};
1252 template <typename T = vec> EnableIfUsingArray<T> operator+() const {
1254 for (size_t I = 0; I < NumElements; ++I)
1255 Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1260 template <typename T = vec> EnableIfNotUsingArray<T> operator-() const {
1262 if constexpr (std::is_same_v<Type, bool>) {
1263 Ret.ConvertToDataT();
1268 template <typename T = vec> EnableIfUsingArray<T> operator-() const {
1270 for (size_t I = 0; I < NumElements; ++I)
1271 Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
1276 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1280 template <typename T = DataT>
1281 typename std::enable_if_t<std::is_integral_v<vec_data_t<T>>, vec>
1285 #ifdef __SYCL_DEVICE_ONLY__
1286 vec Ret{(typename vec::DataType) ~m_Data};
1287 if constexpr (std::is_same_v<Type, bool>) {
1288 Ret.ConvertToDataT();
1293 for (size_t I = 0; I < NumElements; ++I) {
1294 Ret.setValue(I, ~getValue(I));
1300 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
1301 template <typename T = DataT, int N = NumElements>
1302 typename std::enable_if_t<std::is_same_v<std::byte, T>, vec<T, N>>
1306 #ifdef __SYCL_DEVICE_ONLY__
1307 return vec<T, N>{(typename vec<T, N>::DataType) !m_Data};
1310 for (size_t I = 0; I < N; ++I) {
1311 Ret.setValue(I, std::byte{!vec_data<DataT>::get(getValue(I))});
1317 template <typename T = DataT, int N = NumElements>
1318 typename std::enable_if_t<!std::is_same_v<std::byte, T>, vec<T, N>>
1322 #ifdef __SYCL_DEVICE_ONLY__
1323 return vec<T, N>{(typename vec<T, N>::DataType) !m_Data};
1326 for (size_t I = 0; I < N; ++I) {
1327 Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1333 template <typename T = DataT, int N = NumElements>
1334 vec<T, N> operator!() const {
1335 #ifdef __SYCL_DEVICE_ONLY__
1336 return vec<T, N>{(typename vec<T, N>::DataType) !m_Data};
1339 for (size_t I = 0; I < N; ++I) {
1340 Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1347 vec operator+() const {
1350 #ifdef __SYCL_DEVICE_ONLY__
1351 return vec{+m_Data};
1354 for (size_t I = 0; I < NumElements; ++I)
1355 Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1360 vec operator-() const {
1363 #ifdef __SYCL_DEVICE_ONLY__
1365 if constexpr (std::is_same_v<Type, bool>) {
1366 Ret.ConvertToDataT();
1371 for (size_t I = 0; I < NumElements; ++I)
1372 Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
1391 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1392 template <template <typename> class Operation,
1393 typename Ty = vec<DataT, NumElements>>
1394 vec<DataT, NumElements>
1395 operatorHelper(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) const {
1396 vec<DataT, NumElements> Result;
1397 Operation<DataType> Op;
1398 Result.m_Data = Op(m_Data, Rhs.m_Data);
1402 template <template <typename> class Operation,
1403 typename Ty = vec<DataT, NumElements>>
1404 vec<DataT, NumElements>
1405 operatorHelper(const EnableIfUsingArrayOnDevice<Ty> &Rhs) const {
1406 vec<DataT, NumElements> Result;
1407 Operation<DataT> Op;
1408 for (size_t I = 0; I < NumElements; ++I) {
1409 Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1414 template <template <typename> class Operation>
1415 vec<DataT, NumElements>
1416 operatorHelper(const vec<DataT, NumElements> &Rhs) const {
1417 vec<DataT, NumElements> Result;
1418 Operation<DataT> Op;
1419 for (size_t I = 0; I < NumElements; ++I) {
1420 Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1429 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1430 template <int Num = NumElements, typename Ty = int,
1431 typename = typename std::enable_if_t<1 != Num>>
1432 constexpr void setValue(EnableIfNotHostHalf<Ty> Index, const DataT &Value,
1434 m_Data[Index] = vec_data<DataT>::get(Value);
1437 template <int Num = NumElements, typename Ty = int,
1438 typename = typename std::enable_if_t<1 != Num>>
1439 constexpr DataT getValue(EnableIfNotHostHalf<Ty> Index, int) const {
1440 return vec_data<DataT>::get(m_Data[Index]);
1443 template <int Num = NumElements, typename Ty = int,
1444 typename = typename std::enable_if_t<1 != Num>>
1445 constexpr void setValue(EnableIfHostHalf<Ty> Index, const DataT &Value, int) {
1446 m_Data.s[Index] = vec_data<DataT>::get(Value);
1449 template <int Num = NumElements, typename Ty = int,
1450 typename = typename std::enable_if_t<1 != Num>>
1451 constexpr DataT getValue(EnableIfHostHalf<Ty> Index, int) const {
1452 return vec_data<DataT>::get(m_Data.s[Index]);
1455 template <int Num = NumElements,
1456 typename = typename std::enable_if_t<1 != Num>>
1457 constexpr void setValue(int Index, const DataT &Value, int) {
1458 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1459 m_Data[Index] = vec_data<DataT>::get(Value);
1461 m_Data.s[Index] = vec_data<DataT>::get(Value);
1465 template <int Num = NumElements,
1466 typename = typename std::enable_if_t<1 != Num>>
1467 constexpr DataT getValue(int Index, int) const {
1468 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1469 return vec_data<DataT>::get(m_Data[Index]);
1471 return vec_data<DataT>::get(m_Data.s[Index]);
1476 template <int Num = NumElements,
1477 typename = typename std::enable_if_t<1 == Num>>
1478 constexpr void setValue(int, const DataT &Value, float) {
1479 m_Data = vec_data<DataT>::get(Value);
1482 template <int Num = NumElements,
1483 typename = typename std::enable_if_t<1 == Num>>
1484 DataT getValue(int, float) const {
1485 return vec_data<DataT>::get(m_Data);
1489 constexpr void setValue(int Index, const DataT &Value) {
1490 if (NumElements == 1)
1491 setValue(Index, Value, 0);
1493 setValue(Index, Value, 0.f);
1496 DataT getValue(int Index) const {
1497 return (NumElements == 1) ? getValue(Index, 0) : getValue(Index, 0.f);
1500 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1506 alignas(detail::vector_alignment<DataT, NumElements>::value) DataType m_Data;
1510 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
1519 __SYCL_ALIGNED_VAR(DataType,
1520 (detail::vector_alignment<DataT, NumElements>::value),
1526 template <typename T1, typename T2, typename T3, template <typename> class T4,
1528 friend class detail::SwizzleOp;
1529 template <typename T1, int T2> friend class vec;
1532 #ifdef __cpp_deduction_guides
1534 template <class T, class... U,
1535 class = std::enable_if_t<(std::is_same_v<T, U> && ...)>>
1536 vec(T, U...) -> vec<T, sizeof...(U) + 1>;
1543 template <typename VecT, typename OperationLeftT, typename OperationRightT,
1544 template <typename> class OperationCurrentT, int... Indexes>
1546 using DataT = typename VecT::element_type;
1573 using OpLeftDataT = typename OperationLeftT::DataT;
1574 using OpRightDataT = typename OperationRightT::DataT;
1575 using CommonDataT = std::conditional_t<
1576 sizeof(DataT) >= sizeof(std::common_type_t<OpLeftDataT, OpRightDataT>),
1577 DataT, std::common_type_t<OpLeftDataT, OpRightDataT>>;
1578 static constexpr int getNumElements() { return sizeof...(Indexes); }
1580 using rel_t = detail::rel_t<DataT>;
1581 using vec_t = vec<DataT, sizeof...(Indexes)>;
1582 using vec_rel_t = vec<rel_t, sizeof...(Indexes)>;
1584 template <typename OperationRightT_,
1585 template <typename> class OperationCurrentT_, int... Idx_>
1586 using NewLHOp = SwizzleOp<VecT,
1587 SwizzleOp<VecT, OperationLeftT, OperationRightT,
1588 OperationCurrentT, Indexes...>,
1589 OperationRightT_, OperationCurrentT_, Idx_...>;
1591 template <typename OperationRightT_,
1592 template <typename> class OperationCurrentT_, int... Idx_>
1593 using NewRelOp = SwizzleOp<vec<rel_t, VecT::getNumElements()>,
1594 SwizzleOp<VecT, OperationLeftT, OperationRightT,
1595 OperationCurrentT, Indexes...>,
1596 OperationRightT_, OperationCurrentT_, Idx_...>;
1598 template <typename OperationLeftT_,
1599 template <typename> class OperationCurrentT_, int... Idx_>
1600 using NewRHOp = SwizzleOp<VecT, OperationLeftT_,
1601 SwizzleOp<VecT, OperationLeftT, OperationRightT,
1602 OperationCurrentT, Indexes...>,
1603 OperationCurrentT_, Idx_...>;
1605 template <int IdxNum, typename T = void>
1606 using EnableIfOneIndex = typename std::enable_if_t<
1607 1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1609 template <int IdxNum, typename T = void>
1610 using EnableIfMultipleIndexes = typename std::enable_if_t<
1611 1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1613 template <typename T>
1614 using EnableIfScalarType = typename std::enable_if_t<
1615 std::is_convertible_v<DataT, T> &&
1616 (std::is_fundamental_v<vec_data_t<T>> ||
1617 std::is_same_v<typename std::remove_const_t<T>, half>)>;
1619 template <typename T>
1620 using EnableIfNoScalarType = typename std::enable_if_t<
1621 !std::is_convertible_v<DataT, T> ||
1622 !(std::is_fundamental_v<vec_data_t<T>> ||
1623 std::is_same_v<typename std::remove_const_t<T>, half>)>;
1625 template <int... Indices>
1627 SwizzleOp<VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1629 template <int... Indices>
1630 using ConstSwizzle =
1631 SwizzleOp<const VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1634 using element_type = DataT;
1636 const DataT &operator[](int i) const {
1637 std::array<int, getNumElements()> Idxs{Indexes...};
1638 return (*m_Vector)[Idxs[i]];
1641 template <typename _T = VecT>
1642 std::enable_if_t<!std::is_const_v<_T>, DataT> &operator[](int i) {
1643 std::array<int, getNumElements()> Idxs{Indexes...};
1644 return (*m_Vector)[Idxs[i]];
1647 __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1648 size_t get_count() const { return size(); }
1649 size_t size() const noexcept { return getNumElements(); }
1651 template <int Num = getNumElements()>
1652 __SYCL2020_DEPRECATED(
1653 "get_size() is deprecated, please use byte_size() instead")
1654 size_t get_size() const {
1655 return byte_size<Num>();
1658 template <int Num = getNumElements()> size_t byte_size() const noexcept {
1659 return sizeof(DataT) * (Num == 3 ? 4 : Num);
1662 template <typename T, int IdxNum = getNumElements(),
1663 typename = EnableIfOneIndex<IdxNum>,
1664 typename = EnableIfScalarType<T>>
1665 operator T() const {
1669 template <typename T, typename = EnableIfScalarType<T>>
1670 friend NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>
1671 operator*(const T &Lhs, const SwizzleOp &Rhs) {
1672 return NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
1673 Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1676 template <typename T, typename = EnableIfScalarType<T>>
1677 friend NewRHOp<GetScalarOp<T>, std::plus, Indexes...>
1678 operator+(const T &Lhs, const SwizzleOp &Rhs) {
1679 return NewRHOp<GetScalarOp<T>, std::plus, Indexes...>(
1680 Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1683 template <typename T, typename = EnableIfScalarType<T>>
1684 friend NewRHOp<GetScalarOp<T>, std::divides, Indexes...>
1685 operator/(const T &Lhs, const SwizzleOp &Rhs) {
1686 return NewRHOp<GetScalarOp<T>, std::divides, Indexes...>(
1687 Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1691 #ifdef __SYCL_OPASSIGN
1692 #error "Undefine __SYCL_OPASSIGN macro."
1694 #define __SYCL_OPASSIGN(OPASSIGN, OP) \
1695 SwizzleOp &operator OPASSIGN(const DataT & Rhs) { \
1696 operatorHelper<OP>(vec_t(Rhs)); \
1699 template <typename RhsOperation> \
1700 SwizzleOp &operator OPASSIGN(const RhsOperation & Rhs) { \
1701 operatorHelper<OP>(Rhs); \
1705 __SYCL_OPASSIGN(+=, std::plus)
1706 __SYCL_OPASSIGN(-=, std::minus)
1707 __SYCL_OPASSIGN(*=, std::multiplies)
1708 __SYCL_OPASSIGN(/=, std::divides)
1709 __SYCL_OPASSIGN(%=, std::modulus)
1710 __SYCL_OPASSIGN(&=, std::bit_and)
1711 __SYCL_OPASSIGN(|=, std::bit_or)
1712 __SYCL_OPASSIGN(^=, std::bit_xor)
1713 __SYCL_OPASSIGN(>>=, RShift)
1714 __SYCL_OPASSIGN(<<=, LShift)
1715 #undef __SYCL_OPASSIGN
1718 #error "Undefine __SYCL_UOP macro"
1720 #define __SYCL_UOP(UOP, OPASSIGN) \
1721 SwizzleOp &operator UOP() { \
1722 *this OPASSIGN static_cast<DataT>(1); \
1725 vec_t operator UOP(int) { \
1726 vec_t Ret = *this; \
1727 *this OPASSIGN static_cast<DataT>(1); \
1735 template <typename T = DataT>
1736 typename std::enable_if_t<std::is_integral_v<vec_data_t<T>>, vec_t>
1742 vec_rel_t operator!() {
1757 template <int IdxNum = getNumElements(),
1758 typename = EnableIfMultipleIndexes<IdxNum>>
1759 SwizzleOp &operator=(const vec<DataT, IdxNum> &Rhs) {
1760 std::array<int, IdxNum> Idxs{Indexes...};
1761 for (size_t I = 0; I < Idxs.size(); ++I) {
1762 m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1767 template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
1768 SwizzleOp &operator=(const DataT &Rhs) {
1769 std::array<int, IdxNum> Idxs{Indexes...};
1770 m_Vector->setValue(Idxs[0], Rhs);
1774 template <int IdxNum = getNumElements(),
1775 EnableIfMultipleIndexes<IdxNum, bool> = true>
1776 SwizzleOp &operator=(const DataT &Rhs) {
1777 std::array<int, IdxNum> Idxs{Indexes...};
1778 for (auto Idx : Idxs) {
1779 m_Vector->setValue(Idx, Rhs);
1784 template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
1785 SwizzleOp &operator=(DataT &&Rhs) {
1786 std::array<int, IdxNum> Idxs{Indexes...};
1787 m_Vector->setValue(Idxs[0], Rhs);
1791 template <typename T, typename = EnableIfScalarType<T>>
1792 NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>
1793 operator*(const T &Rhs) const {
1794 return NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
1795 m_Vector, *this, GetScalarOp<T>(Rhs));
1798 template <typename RhsOperation,
1799 typename = EnableIfNoScalarType<RhsOperation>>
1800 NewLHOp<RhsOperation, std::multiplies, Indexes...>
1801 operator*(const RhsOperation &Rhs) const {
1802 return NewLHOp<RhsOperation, std::multiplies, Indexes...>(m_Vector, *this,
1806 template <typename T, typename = EnableIfScalarType<T>>
1807 NewLHOp<GetScalarOp<T>, std::plus, Indexes...> operator+(const T &Rhs) const {
1808 return NewLHOp<GetScalarOp<T>, std::plus, Indexes...>(m_Vector, *this,
1809 GetScalarOp<T>(Rhs));
1812 template <typename RhsOperation,
1813 typename = EnableIfNoScalarType<RhsOperation>>
1814 NewLHOp<RhsOperation, std::plus, Indexes...>
1815 operator+(const RhsOperation &Rhs) const {
1816 return NewLHOp<RhsOperation, std::plus, Indexes...>(m_Vector, *this, Rhs);
1819 template <typename T, typename = EnableIfScalarType<T>>
1820 NewLHOp<GetScalarOp<T>, std::minus, Indexes...>
1821 operator-(const T &Rhs) const {
1822 return NewLHOp<GetScalarOp<T>, std::minus, Indexes...>(m_Vector, *this,
1823 GetScalarOp<T>(Rhs));
1826 template <typename RhsOperation,
1827 typename = EnableIfNoScalarType<RhsOperation>>
1828 NewLHOp<RhsOperation, std::minus, Indexes...>
1829 operator-(const RhsOperation &Rhs) const {
1830 return NewLHOp<RhsOperation, std::minus, Indexes...>(m_Vector, *this, Rhs);
1833 template <typename T, typename = EnableIfScalarType<T>>
1834 NewLHOp<GetScalarOp<T>, std::divides, Indexes...>
1835 operator/(const T &Rhs) const {
1836 return NewLHOp<GetScalarOp<T>, std::divides, Indexes...>(
1837 m_Vector, *this, GetScalarOp<T>(Rhs));
1840 template <typename RhsOperation,
1841 typename = EnableIfNoScalarType<RhsOperation>>
1842 NewLHOp<RhsOperation, std::divides, Indexes...>
1843 operator/(const RhsOperation &Rhs) const {
1844 return NewLHOp<RhsOperation, std::divides, Indexes...>(m_Vector, *this,
1848 template <typename T, typename = EnableIfScalarType<T>>
1849 NewLHOp<GetScalarOp<T>, std::modulus, Indexes...>
1850 operator%(const T &Rhs) const {
1851 return NewLHOp<GetScalarOp<T>, std::modulus, Indexes...>(
1852 m_Vector, *this, GetScalarOp<T>(Rhs));
1855 template <typename RhsOperation,
1856 typename = EnableIfNoScalarType<RhsOperation>>
1857 NewLHOp<RhsOperation, std::modulus, Indexes...>
1858 operator%(const RhsOperation &Rhs) const {
1859 return NewLHOp<RhsOperation, std::modulus, Indexes...>(m_Vector, *this,
1863 template <typename T, typename = EnableIfScalarType<T>>
1864 NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>
1865 operator&(const T &Rhs) const {
1866 return NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>(
1867 m_Vector, *this, GetScalarOp<T>(Rhs));
1870 template <typename RhsOperation,
1871 typename = EnableIfNoScalarType<RhsOperation>>
1872 NewLHOp<RhsOperation, std::bit_and, Indexes...>
1873 operator&(const RhsOperation &Rhs) const {
1874 return NewLHOp<RhsOperation, std::bit_and, Indexes...>(m_Vector, *this,
1878 template <typename T, typename = EnableIfScalarType<T>>
1879 NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>
1880 operator|(const T &Rhs) const {
1881 return NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>(
1882 m_Vector, *this, GetScalarOp<T>(Rhs));
1885 template <typename RhsOperation,
1886 typename = EnableIfNoScalarType<RhsOperation>>
1887 NewLHOp<RhsOperation, std::bit_or, Indexes...>
1888 operator|(const RhsOperation &Rhs) const {
1889 return NewLHOp<RhsOperation, std::bit_or, Indexes...>(m_Vector, *this, Rhs);
1892 template <typename T, typename = EnableIfScalarType<T>>
1893 NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>
1894 operator^(const T &Rhs) const {
1895 return NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>(
1896 m_Vector, *this, GetScalarOp<T>(Rhs));
1899 template <typename RhsOperation,
1900 typename = EnableIfNoScalarType<RhsOperation>>
1901 NewLHOp<RhsOperation, std::bit_xor, Indexes...>
1902 operator^(const RhsOperation &Rhs) const {
1903 return NewLHOp<RhsOperation, std::bit_xor, Indexes...>(m_Vector, *this,
1907 template <typename T, typename = EnableIfScalarType<T>>
1908 NewLHOp<GetScalarOp<T>, RShift, Indexes...> operator>>(const T &Rhs) const {
1909 return NewLHOp<GetScalarOp<T>, RShift, Indexes...>(m_Vector, *this,
1910 GetScalarOp<T>(Rhs));
1913 template <typename RhsOperation,
1914 typename = EnableIfNoScalarType<RhsOperation>>
1915 NewLHOp<RhsOperation, RShift, Indexes...>
1916 operator>>(const RhsOperation &Rhs) const {
1917 return NewLHOp<RhsOperation, RShift, Indexes...>(m_Vector, *this, Rhs);
1920 template <typename T, typename = EnableIfScalarType<T>>
1921 NewLHOp<GetScalarOp<T>, LShift, Indexes...> operator<<(const T &Rhs) const {
1922 return NewLHOp<GetScalarOp<T>, LShift, Indexes...>(m_Vector, *this,
1923 GetScalarOp<T>(Rhs));
1926 template <typename RhsOperation,
1927 typename = EnableIfNoScalarType<RhsOperation>>
1928 NewLHOp<RhsOperation, LShift, Indexes...>
1929 operator<<(const RhsOperation &Rhs) const {
1930 return NewLHOp<RhsOperation, LShift, Indexes...>(m_Vector, *this, Rhs);
1934 typename T1, typename T2, typename T3, template <typename> class T4,
1936 typename = typename std::enable_if_t<sizeof...(T5) == getNumElements()>>
1937 SwizzleOp &operator=(const SwizzleOp<T1, T2, T3, T4, T5...> &Rhs) {
1938 std::array<int, getNumElements()> Idxs{Indexes...};
1939 for (size_t I = 0; I < Idxs.size(); ++I) {
1940 m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1946 typename T1, typename T2, typename T3, template <typename> class T4,
1948 typename = typename std::enable_if_t<sizeof...(T5) == getNumElements()>>
1949 SwizzleOp &operator=(SwizzleOp<T1, T2, T3, T4, T5...> &&Rhs) {
1950 std::array<int, getNumElements()> Idxs{Indexes...};
1951 for (size_t I = 0; I < Idxs.size(); ++I) {
1952 m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1957 template <typename T, typename = EnableIfScalarType<T>>
1958 NewRelOp<GetScalarOp<T>, EqualTo, Indexes...> operator==(const T &Rhs) const {
1959 return NewRelOp<GetScalarOp<T>, EqualTo, Indexes...>(NULL, *this,
1960 GetScalarOp<T>(Rhs));
1963 template <typename RhsOperation,
1964 typename = EnableIfNoScalarType<RhsOperation>>
1965 NewRelOp<RhsOperation, EqualTo, Indexes...>
1966 operator==(const RhsOperation &Rhs) const {
1967 return NewRelOp<RhsOperation, EqualTo, Indexes...>(NULL, *this, Rhs);
1970 template <typename T, typename = EnableIfScalarType<T>>
1971 NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>
1972 operator!=(const T &Rhs) const {
1973 return NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>(
1974 NULL, *this, GetScalarOp<T>(Rhs));
1977 template <typename RhsOperation,
1978 typename = EnableIfNoScalarType<RhsOperation>>
1979 NewRelOp<RhsOperation, NotEqualTo, Indexes...>
1980 operator!=(const RhsOperation &Rhs) const {
1981 return NewRelOp<RhsOperation, NotEqualTo, Indexes...>(NULL, *this, Rhs);
1984 template <typename T, typename = EnableIfScalarType<T>>
1985 NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>
1986 operator>=(const T &Rhs) const {
1987 return NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>(
1988 NULL, *this, GetScalarOp<T>(Rhs));
1991 template <typename RhsOperation,
1992 typename = EnableIfNoScalarType<RhsOperation>>
1993 NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>
1994 operator>=(const RhsOperation &Rhs) const {
1995 return NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>(NULL, *this, Rhs);
1998 template <typename T, typename = EnableIfScalarType<T>>
1999 NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>
2000 operator<=(const T &Rhs) const {
2001 return NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>(
2002 NULL, *this, GetScalarOp<T>(Rhs));
2005 template <typename RhsOperation,
2006 typename = EnableIfNoScalarType<RhsOperation>>
2007 NewRelOp<RhsOperation, LessEqualTo, Indexes...>
2008 operator<=(const RhsOperation &Rhs) const {
2009 return NewRelOp<RhsOperation, LessEqualTo, Indexes...>(NULL, *this, Rhs);
2012 template <typename T, typename = EnableIfScalarType<T>>
2013 NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>
2014 operator>(const T &Rhs) const {
2015 return NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>(
2016 NULL, *this, GetScalarOp<T>(Rhs));
2019 template <typename RhsOperation,
2020 typename = EnableIfNoScalarType<RhsOperation>>
2021 NewRelOp<RhsOperation, GreaterThan, Indexes...>
2022 operator>(const RhsOperation &Rhs) const {
2023 return NewRelOp<RhsOperation, GreaterThan, Indexes...>(NULL, *this, Rhs);
2026 template <typename T, typename = EnableIfScalarType<T>>
2027 NewRelOp<GetScalarOp<T>, LessThan, Indexes...> operator<(const T &Rhs) const {
2028 return NewRelOp<GetScalarOp<T>, LessThan, Indexes...>(NULL, *this,
2029 GetScalarOp<T>(Rhs));
2032 template <typename RhsOperation,
2033 typename = EnableIfNoScalarType<RhsOperation>>
2034 NewRelOp<RhsOperation, LessThan, Indexes...>
2035 operator<(const RhsOperation &Rhs) const {
2036 return NewRelOp<RhsOperation, LessThan, Indexes...>(NULL, *this, Rhs);
2039 template <typename T, typename = EnableIfScalarType<T>>
2040 NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>
2041 operator&&(const T &Rhs) const {
2042 return NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>(
2043 NULL, *this, GetScalarOp<T>(Rhs));
2046 template <typename RhsOperation,
2047 typename = EnableIfNoScalarType<RhsOperation>>
2048 NewRelOp<RhsOperation, LogicalAnd, Indexes...>
2049 operator&&(const RhsOperation &Rhs) const {
2050 return NewRelOp<RhsOperation, LogicalAnd, Indexes...>(NULL, *this, Rhs);
2053 template <typename T, typename = EnableIfScalarType<T>>
2054 NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>
2055 operator||(const T &Rhs) const {
2056 return NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>(NULL, *this,
2057 GetScalarOp<T>(Rhs));
2060 template <typename RhsOperation,
2061 typename = EnableIfNoScalarType<RhsOperation>>
2062 NewRelOp<RhsOperation, LogicalOr, Indexes...>
2063 operator||(const RhsOperation &Rhs) const {
2064 return NewRelOp<RhsOperation, LogicalOr, Indexes...>(NULL, *this, Rhs);
2074 template <int Index> struct Indexer {
2075 static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...};
2076 static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index];
2080 #ifdef __SYCL_ACCESS_RETURN
2081 #error "Undefine __SYCL_ACCESS_RETURN macro"
2083 #define __SYCL_ACCESS_RETURN m_Vector
2084 #include "swizzles.def"
2085 #undef __SYCL_ACCESS_RETURN
2090 template <access::address_space Space, access::decorated DecorateAddress>
2091 void load(size_t offset, multi_ptr<DataT, Space, DecorateAddress> ptr) {
2093 Tmp.template load(offset, ptr);
2097 template <typename convertT, rounding_mode roundingMode>
2098 vec<convertT, sizeof...(Indexes)> convert() const {
2101 return Tmp.template convert<convertT, roundingMode>();
2104 template <typename asT> asT as() const {
2107 static_assert((sizeof(Tmp) == sizeof(asT)),
2108 "The new SYCL vec type must have the same storage size in "
2109 "bytes as this SYCL swizzled vec");
2111 detail::is_contained<asT, detail::gtl::vector_basic_list>::value ||
2112 detail::is_contained<asT, detail::gtl::vector_bool_list>::value,
2113 "asT must be SYCL vec of a different element type and "
2114 "number of elements specified by asT");
2115 return Tmp.template as<asT>();
2119 SwizzleOp(const SwizzleOp &Rhs)
2120 : m_Vector(Rhs.m_Vector), m_LeftOperation(Rhs.m_LeftOperation),
2121 m_RightOperation(Rhs.m_RightOperation) {}
2123 SwizzleOp(VecT *Vector, OperationLeftT LeftOperation,
2124 OperationRightT RightOperation)
2125 : m_Vector(Vector), m_LeftOperation(LeftOperation),
2126 m_RightOperation(RightOperation) {}
2128 SwizzleOp(VecT *Vector) : m_Vector(Vector) {}
2130 SwizzleOp(SwizzleOp &&Rhs)
2131 : m_Vector(Rhs.m_Vector), m_LeftOperation(std::move(Rhs.m_LeftOperation)),
2132 m_RightOperation(std::move(Rhs.m_RightOperation)) {}
2138 template <int IdxNum = getNumElements()>
2139 CommonDataT getValue(EnableIfOneIndex<IdxNum, size_t> Index) const {
2140 if (std::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
2141 std::array<int, getNumElements()> Idxs{Indexes...};
2142 return m_Vector->getValue(Idxs[Index]);
2144 auto Op = OperationCurrentT<vec_data_t<CommonDataT>>();
2145 return vec_data<CommonDataT>::get(
2146 Op(vec_data<CommonDataT>::get(m_LeftOperation.getValue(Index)),
2147 vec_data<CommonDataT>::get(m_RightOperation.getValue(Index))));
2150 template <int IdxNum = getNumElements()>
2151 DataT getValue(EnableIfMultipleIndexes<IdxNum, size_t> Index) const {
2152 if (std::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
2153 std::array<int, getNumElements()> Idxs{Indexes...};
2154 return m_Vector->getValue(Idxs[Index]);
2156 auto Op = OperationCurrentT<vec_data_t<DataT>>();
2157 return vec_data<DataT>::get(
2158 Op(vec_data<DataT>::get(m_LeftOperation.getValue(Index)),
2159 vec_data<DataT>::get(m_RightOperation.getValue(Index))));
2162 template <template <typename> class Operation, typename RhsOperation>
2163 void operatorHelper(const RhsOperation &Rhs) {
2164 Operation<vec_data_t<DataT>> Op;
2165 std::array<int, getNumElements()> Idxs{Indexes...};
2166 for (size_t I = 0; I < Idxs.size(); ++I) {
2167 DataT Res = vec_data<DataT>::get(
2168 Op(vec_data<DataT>::get(m_Vector->getValue(Idxs[I])),
2169 vec_data<DataT>::get(Rhs.getValue(I))));
2170 m_Vector->setValue(Idxs[I], Res);
2177 OperationLeftT m_LeftOperation;
2178 OperationRightT m_RightOperation;
2181 template <typename T1, int T2> friend class sycl::vec;
2183 template <typename T1, typename T2, typename T3, template <typename> class T4,
2185 friend class SwizzleOp;
2193 #error "Undefine __SYCL_BINOP macro"
2195 #define __SYCL_BINOP(BINOP) \
2196 template <typename T, int Num> \
2197 typename std::enable_if_t< \
2198 std::is_fundamental_v<vec_data_t<T>> || \
2199 std::is_same_v<typename std::remove_const_t<T>, half>, \
2201 operator BINOP(const T & Lhs, const vec<T, Num> &Rhs) { \
2202 return vec<T, Num>(Lhs) BINOP Rhs; \
2204 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2205 template <typename> class OperationCurrentT, int... Indexes, \
2206 typename T, typename T1 = typename VecT::element_type, \
2207 int Num = sizeof...(Indexes)> \
2208 typename std::enable_if_t< \
2209 std::is_convertible_v<T, T1> && \
2210 (std::is_fundamental_v<vec_data_t<T>> || \
2211 std::is_same_v<typename std::remove_const_t<T>, half>), \
2215 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2216 OperationCurrentT, Indexes...> &Rhs) { \
2217 vec<T1, Num> Tmp = Rhs; \
2218 return Lhs BINOP Tmp; \
2220 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2221 template <typename> class OperationCurrentT, int... Indexes, \
2222 typename T = typename VecT::element_type, \
2223 int Num = sizeof...(Indexes)> \
2224 vec<T, Num> operator BINOP( \
2225 const vec<T, Num> &Lhs, \
2226 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2227 OperationCurrentT, Indexes...> &Rhs) { \
2228 vec<T, Num> Tmp = Rhs; \
2229 return Lhs BINOP Tmp; \
2247 #ifdef __SYCL_RELLOGOP
2248 #error "Undefine __SYCL_RELLOGOP macro"
2250 #define __SYCL_RELLOGOP(RELLOGOP) \
2251 template <typename T, typename DataT, int Num> \
2252 typename std::enable_if_t< \
2253 std::is_convertible_v<T, DataT> && \
2254 (std::is_fundamental_v<vec_data_t<T>> || \
2255 std::is_same_v<typename std::remove_const_t<T>, half>), \
2256 vec<detail::rel_t<DataT>, Num>> \
2257 operator RELLOGOP(const T & Lhs, const vec<DataT, Num> &Rhs) { \
2258 return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs; \
2260 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2261 template <typename> class OperationCurrentT, int... Indexes, \
2262 typename T, typename T1 = typename VecT::element_type, \
2263 int Num = sizeof...(Indexes)> \
2264 typename std::enable_if_t< \
2265 std::is_convertible_v<T, T1> && \
2266 (std::is_fundamental_v<vec_data_t<T>> || \
2267 std::is_same_v<typename std::remove_const_t<T>, half>), \
2268 vec<detail::rel_t<T1>, Num>> \
2269 operator RELLOGOP( \
2271 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2272 OperationCurrentT, Indexes...> &Rhs) { \
2273 vec<T1, Num> Tmp = Rhs; \
2274 return Lhs RELLOGOP Tmp; \
2276 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2277 template <typename> class OperationCurrentT, int... Indexes, \
2278 typename T = typename VecT::element_type, \
2279 int Num = sizeof...(Indexes)> \
2280 vec<detail::rel_t<T>, Num> operator RELLOGOP( \
2281 const vec<T, Num> &Lhs, \
2282 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2283 OperationCurrentT, Indexes...> &Rhs) { \
2284 vec<T, Num> Tmp = Rhs; \
2285 return Lhs RELLOGOP Tmp; \
2297 #undef __SYCL_RELLOGOP
2303 constexpr bool isValidVectorSize(int N) {
2304 return N == 2 || N == 3 || N == 4 || N == 8 || N == 16;
2306 template <typename T, int N, typename V> struct VecStorage {
2308 isValidVectorSize(N) || N == 1,
2309 "Incorrect number of elements for sycl::vec: only 1, 2, 3, 4, 8 "
2310 "or 16 are supported");
2311 static_assert(!std::is_same_v<V, void>, "Incorrect data type for sycl::vec");
2314 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2316 #ifdef __SYCL_DEVICE_ONLY__
2320 template <typename T, int N> struct VecStorageImpl {
2321 static constexpr size_t Num = (N == 3) ? 4 : N;
2322 static constexpr size_t Sz = Num * sizeof(T);
2324 typename std::conditional<Sz <= 64, T __attribute__((ext_vector_type(N))),
2325 std::array<T, Num>>::type;
2326 using VectorDataType = T __attribute__((ext_vector_type(N)));
2329 template <typename T, int N> struct VecStorageImpl {
2330 using DataType = std::array<T, (N == 3) ? 4 : N>;
2335 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2337 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
2338 template <typename T, int N> struct VecStorageImpl {
2339 using DataType = T __attribute__((ext_vector_type(N)));
2344 template <typename T, int N> struct VecStorageImpl;
2345 #define __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, num) \
2346 template <> struct VecStorageImpl<type, num> { \
2347 using DataType = ::cl_##cl_type##num; \
2351 #ifndef __SYCL_USE_EXT_VECTOR_TYPE__
2352 #define __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(type, cl_type) \
2353 __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 2) \
2354 __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 3) \
2355 __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 4) \
2356 __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 8) \
2357 __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, 16)
2359 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int8_t, char)
2360 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int16_t, short)
2361 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int32_t, int)
2362 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::int64_t, long)
2363 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint8_t, uchar)
2364 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint16_t, ushort)
2365 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint32_t, uint)
2366 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(std::uint64_t, ulong)
2367 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(float, float)
2368 __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(double, double)
2370 #undef __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE
2371 #undef __SYCL_DEFINE_VECSTORAGE_IMPL
2376 template <> struct VecStorage<bool, 1, void> {
2377 using DataType = bool;
2378 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2379 #ifdef __SYCL_DEVICE_ONLY__
2380 using VectorDataType = bool;
2387 struct VecStorage<bool, N, typename std::enable_if_t<isValidVectorSize(N)>> {
2389 typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
2390 std::int32_t, std::int64_t>,
2392 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2393 #ifdef __SYCL_DEVICE_ONLY__
2394 using VectorDataType =
2395 typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
2396 std::int32_t, std::int64_t>,
2402 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2405 template <> struct VecStorage<std::byte, 1, void> {
2406 using DataType = std::int8_t;
2407 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2408 #ifdef __SYCL_DEVICE_ONLY__
2409 using VectorDataType = std::int8_t;
2416 template <typename T>
2417 struct VecStorage<T, 1, typename std::enable_if_t<is_sigeninteger_v<T>>> {
2418 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
2419 using DataType = select_apply_cl_t<T, std::int8_t, std::int16_t, std::int32_t,
2423 #ifdef __SYCL_DEVICE_ONLY__
2424 using VectorDataType = DataType;
2430 template <typename T>
2431 struct VecStorage<T, 1, typename std::enable_if_t<is_sugeninteger_v<T>>> {
2432 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
2433 using DataType = select_apply_cl_t<T, std::uint8_t, std::uint16_t,
2434 std::uint32_t, std::uint64_t>;
2437 #ifdef __SYCL_DEVICE_ONLY__
2438 using VectorDataType = DataType;
2444 template <typename T>
2446 T, 1, typename std::enable_if_t<!is_half_v<T> && is_sgenfloat_v<T>>> {
2448 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2449 #ifdef __SYCL_DEVICE_ONLY__
2450 using VectorDataType = DataType;
2455 template <typename T, int N>
2458 typename std::enable_if_t<isValidVectorSize(N) &&
2459 (is_sgeninteger_v<T> ||
2460 (is_sgenfloat_v<T> && !is_half_v<T>))>> {
2462 typename VecStorageImpl<typename VecStorage<T, 1>::DataType, N>::DataType;
2463 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2464 #ifdef __SYCL_DEVICE_ONLY__
2465 using VectorDataType =
2466 typename VecStorageImpl<typename VecStorage<T, 1>::DataType,
2473 template <> struct VecStorage<half, 1, void> {
2474 using DataType = sycl::detail::half_impl::StorageT;
2475 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2476 #ifdef __SYCL_DEVICE_ONLY__
2477 using VectorDataType = sycl::detail::half_impl::StorageT;
2483 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
2484 #define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
2485 template <> struct VecStorage<half, Num, void> { \
2486 using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
2487 using VectorDataType = sycl::detail::half_impl::Vec##Num##StorageT; \
2491 #define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
2492 template <> struct VecStorage<half, Num, void> { \
2493 using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
2498 __SYCL_DEFINE_HALF_VECSTORAGE(2)
2499 __SYCL_DEFINE_HALF_VECSTORAGE(3)
2500 __SYCL_DEFINE_HALF_VECSTORAGE(4)
2501 __SYCL_DEFINE_HALF_VECSTORAGE(8)
2502 __SYCL_DEFINE_HALF_VECSTORAGE(16)
2503 #undef __SYCL_DEFINE_HALF_VECSTORAGE
2509 #define SYCL_DEVICE_COPYABLE 1
2518 template <typename T> struct is_device_copyable;
2521 template <typename T, typename = void>
2522 struct is_device_copyable_impl : std::is_trivially_copyable<T> {};
2524 template <typename T>
2525 struct is_device_copyable_impl<
2526 T, std::enable_if_t<!std::is_same_v<T, std::remove_cv_t<T>>>>
2529 : is_device_copyable<std::remove_cv_t<T>> {};
2532 template <typename T>
2533 struct is_device_copyable : detail::is_device_copyable_impl<T> {};
2536 template <typename T>
2537 struct is_device_copyable<std::array<T, 0>> : std::true_type {};
2540 template <typename T, std::size_t N>
2541 struct is_device_copyable<std::array<T, N>> : is_device_copyable<T> {};
2544 template <typename T>
2545 struct is_device_copyable<std::optional<T>> : is_device_copyable<T> {};
2549 template <typename T1, typename T2>
2550 struct is_device_copyable<std::pair<T1, T2>>
2551 : std::bool_constant<is_device_copyable<T1>::value &&
2552 is_device_copyable<T2>::value> {};
2556 template <typename... Ts>
2557 struct is_device_copyable<std::tuple<Ts...>>
2558 : std::bool_constant<(... && is_device_copyable<Ts>::value)> {};
2562 template <typename... Ts>
2563 struct is_device_copyable<std::variant<Ts...>>
2564 : std::bool_constant<(... && is_device_copyable<Ts>::value)> {};
2567 template <typename T, std::size_t N>
2568 struct is_device_copyable<sycl::marray<T, N>> : is_device_copyable<T> {};
2571 template <typename T, std::size_t N>
2572 struct is_device_copyable<T[N]> : is_device_copyable<T> {};
2574 template <typename T>
2575 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
2578 template <typename T, typename = void>
2579 struct IsDeprecatedDeviceCopyable : std::false_type {};
2583 template <typename T>
2584 struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
2585 IsDeprecatedDeviceCopyable<
2586 T, std::enable_if_t<std::is_trivially_copy_constructible_v<T> &&
2587 std::is_trivially_destructible_v<T> &&
2588 !is_device_copyable_v<T>>> : std::true_type {};
2590 template <typename T, int N>
2591 struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
2592 IsDeprecatedDeviceCopyable<T[N]> : IsDeprecatedDeviceCopyable<T> {};
2594 #ifdef __SYCL_DEVICE_ONLY__
2597 template <typename T, unsigned NumFieldsToCheck>
2598 struct CheckFieldsAreDeviceCopyable
2599 : CheckFieldsAreDeviceCopyable<T, NumFieldsToCheck - 1> {
2600 using FieldT = decltype(__builtin_field_type(T, NumFieldsToCheck - 1));
2601 static_assert(is_device_copyable_v<FieldT> ||
2602 detail::IsDeprecatedDeviceCopyable<FieldT>::value,
2603 "The specified type is not device copyable");
2606 template <typename T> struct CheckFieldsAreDeviceCopyable<T, 0> {};
2610 template <typename T, unsigned NumBasesToCheck>
2611 struct CheckBasesAreDeviceCopyable
2612 : CheckBasesAreDeviceCopyable<T, NumBasesToCheck - 1> {
2613 using BaseT = decltype(__builtin_base_type(T, NumBasesToCheck - 1));
2614 static_assert(is_device_copyable_v<BaseT> ||
2615 detail::IsDeprecatedDeviceCopyable<BaseT>::value,
2616 "The specified type is not device copyable");
2619 template <typename T> struct CheckBasesAreDeviceCopyable<T, 0> {};
2632 template <typename FuncT>
2633 struct CheckDeviceCopyable
2634 : CheckFieldsAreDeviceCopyable<FuncT, __builtin_num_fields(FuncT)>,
2635 CheckBasesAreDeviceCopyable<FuncT, __builtin_num_bases(FuncT)> {};
2639 template <typename TransformedArgType, int Dims, typename KernelType>
2640 struct CheckDeviceCopyable<
2641 RoundedRangeKernel<TransformedArgType, Dims, KernelType>>
2642 : CheckDeviceCopyable<KernelType> {};
2644 template <typename TransformedArgType, int Dims, typename KernelType>
2645 struct CheckDeviceCopyable<
2646 RoundedRangeKernelWithKH<TransformedArgType, Dims, KernelType>>
2647 : CheckDeviceCopyable<KernelType> {};
2655 #if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
2656 #undef __SYCL_ALIGNED_VAR