18 #ifndef __NO_EXT_VECTOR_TYPE_ON_HOST__
25 #ifndef __has_extension
26 #define __has_extension(x) 0
28 #ifdef __HAS_EXT_VECTOR_TYPE__
29 #error "Undefine __HAS_EXT_VECTOR_TYPE__ macro"
31 #if __has_extension(attribute_ext_vector_type)
32 #define __HAS_EXT_VECTOR_TYPE__
36 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
37 #error "Undefine __SYCL_USE_EXT_VECTOR_TYPE__ macro"
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__
43 #elif defined(__SYCL_DEVICE_ONLY__)
46 #error "SYCL device compiler is built without ext_vector_type support"
47 #endif // __HAS_EXT_VECTOR_TYPE__
58 #ifndef __SYCL_USE_EXT_VECTOR_TYPE__
67 #ifndef __SYCL_DEVICE_ONLY__
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;
112 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
122 template <
typename VecT,
typename OperationLeftT,
typename OperationRightT,
123 template <
typename>
class OperationCurrentT,
int... Indexes>
126 template <
typename T,
int N,
typename V =
void>
struct VecStorage;
129 template <
typename DataT>
164 return (Lhs == Rhs) ? -1 : 0;
170 return (Lhs != Rhs) ? -1 : 0;
176 return (Lhs >= Rhs) ? -1 : 0;
182 return (Lhs <= Rhs) ? -1 : 0;
188 return (Lhs > Rhs) ? -1 : 0;
194 return (Lhs < Rhs) ? -1 : 0;
200 return (Lhs && Rhs) ? -1 : 0;
206 return (Lhs || Rhs) ? -1 : 0;
222 template <
typename T,
typename R>
224 std::integral_constant<bool, std::is_integral<T>::value &&
225 std::is_integral<R>::value>;
227 template <
typename T,
typename R>
229 std::integral_constant<bool, is_sigeninteger<T>::value &&
232 template <
typename T,
typename R>
234 std::integral_constant<bool, is_sugeninteger<T>::value &&
237 template <
typename T,
typename R>
242 template <
typename T,
typename R>
244 std::integral_constant<bool, std::is_integral<T>::value &&
245 !(std::is_unsigned<T>::value) &&
248 template <
typename T,
typename R>
250 std::integral_constant<bool, std::is_unsigned<T>::value &&
253 template <
typename T,
typename R>
255 std::integral_constant<bool, std::is_integral<T>::value &&
258 template <
typename T,
typename R>
260 std::integral_constant<bool, detail::is_floating_point<T>::value &&
261 std::is_integral<R>::value>;
263 template <
typename T,
typename R>
265 std::integral_constant<bool, detail::is_floating_point<T>::value &&
267 template <
typename T>
269 std::integral_constant<bool, detail::is_sgentype<T>::value>;
271 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
277 #ifndef __SYCL_DEVICE_ONLY__
281 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
284 (is_int_to_int<T, R>::value ||
285 is_int_to_float<T, R>::value ||
286 is_float_to_float<T, R>::value),
289 return static_cast<R
>(Value);
293 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
296 switch (roundingMode) {
298 case rounding_mode::automatic:
300 case rounding_mode::rte: {
301 int OldRoundingDirection = std::fegetround();
302 int Err = std::fesetround(FE_TONEAREST);
304 throw runtime_error(
"Unable to set rounding mode to FE_TONEAREST",
307 Err = std::fesetround(OldRoundingDirection);
309 throw runtime_error(
"Unable to restore rounding mode.", PI_ERROR_UNKNOWN);
313 case rounding_mode::rtz:
316 case rounding_mode::rtp:
319 case rounding_mode::rtn:
322 assert(
false &&
"Unsupported rounding mode!");
323 return static_cast<R
>(Value);
327 template <rounding_mode Mode>
329 Mode == rounding_mode::rte>;
331 template <rounding_mode Mode>
334 template <rounding_mode Mode>
337 template <rounding_mode Mode>
341 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
344 !std::is_same<T, R>::value && std::is_same<OpenCLT, OpenCLR>::value, R>
346 return static_cast<R
>(Value);
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)), \
360 convertImpl(T Value) { \
361 OpenCLT OpValue = sycl::detail::convertDataToType<T, OpenCLT>(Value); \
362 return __spirv_SConvert##_R##DestType(OpValue); \
365 __SYCL_GENERATE_CONVERT_IMPL(
char)
366 __SYCL_GENERATE_CONVERT_IMPL(
short)
367 __SYCL_GENERATE_CONVERT_IMPL(
int)
368 __SYCL_GENERATE_CONVERT_IMPL(
long)
370 #undef __SYCL_GENERATE_CONVERT_IMPL
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, \
380 convertImpl(T Value) { \
381 OpenCLT OpValue = sycl::detail::convertDataToType<T, OpenCLT>(Value); \
382 return __spirv_UConvert##_R##DestType(OpValue); \
385 __SYCL_GENERATE_CONVERT_IMPL(
uchar)
386 __SYCL_GENERATE_CONVERT_IMPL(
ushort)
387 __SYCL_GENERATE_CONVERT_IMPL(
uint)
388 __SYCL_GENERATE_CONVERT_IMPL(ulong)
390 #undef __SYCL_GENERATE_CONVERT_IMPL
393 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
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,
400 return static_cast<R
>(Value);
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)), \
412 convertImpl(T Value) { \
413 OpenCLT OpValue = sycl::detail::convertDataToType<T, OpenCLT>(Value); \
414 return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
417 __SYCL_GENERATE_CONVERT_IMPL(SToF,
half)
418 __SYCL_GENERATE_CONVERT_IMPL(SToF,
float)
419 __SYCL_GENERATE_CONVERT_IMPL(SToF,
double)
421 #undef __SYCL_GENERATE_CONVERT_IMPL
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)), \
432 convertImpl(T Value) { \
433 OpenCLT OpValue = sycl::detail::convertDataToType<T, OpenCLT>(Value); \
434 return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
437 __SYCL_GENERATE_CONVERT_IMPL(UToF,
half)
438 __SYCL_GENERATE_CONVERT_IMPL(UToF,
float)
439 __SYCL_GENERATE_CONVERT_IMPL(UToF,
double)
441 #undef __SYCL_GENERATE_CONVERT_IMPL
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, \
455 convertImpl(T Value) { \
456 OpenCLT OpValue = sycl::detail::convertDataToType<T, OpenCLT>(Value); \
457 return __spirv_FConvert##_R##DestType##_##RoundingMode(OpValue); \
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)
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)
471 #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
472 #undef __SYCL_GENERATE_CONVERT_IMPL
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, \
486 convertImpl(T Value) { \
487 OpenCLT OpValue = sycl::detail::convertDataToType<T, OpenCLT>(Value); \
488 return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue); \
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)
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)
513 #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
514 #undef __SYCL_GENERATE_CONVERT_IMPL
517 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
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,
525 return static_cast<R
>(Value);
528 #endif // __SYCL_DEVICE_ONLY__
531 template <
typename TransformedArgType,
int Dims,
typename KernelType>
532 class RoundedRangeKernel;
533 template <
typename TransformedArgType,
int Dims,
typename KernelType>
534 class RoundedRangeKernelWithKH;
540 template <
typename T>
543 #if defined(_WIN32) && (_MSC_VER)
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
560 #define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
567 template <
typename Type,
int NumElements>
class vec {
572 using DataType =
typename detail::VecStorage<DataT, NumElements>::DataType;
574 static constexpr
int getNumElements() {
return NumElements; }
577 template <
int Counter,
int MaxValue,
class...>
579 std::true_type, std::false_type> {
582 template <
int Counter,
int MaxValue,
typename DataT_,
class... tail>
583 struct SizeChecker<Counter, MaxValue, DataT_, tail...>
585 SizeChecker<Counter + 1, MaxValue, tail...>,
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))...};
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))...};
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))...};
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>());
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)>());
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)>());
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)};
629 template <typename DataT_, typename T> struct FlattenVecArg {
630 constexpr auto operator()(const T &A) const {
631 return FlattenVecArgHelper<DataT_>(A);
636 template <typename DataT_, typename... ArgTN>
637 using VecArgArrayCreator =
638 detail::ArrayCreator<DataT_, FlattenVecArg, ArgTN...>;
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, \
650 struct SizeChecker< \
652 detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
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, \
661 struct SizeChecker< \
663 detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
665 : detail::conditional_t< \
666 Counter + sizeof...(T5) <= MaxValue, \
667 SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
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
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> {};
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 {};
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
708 template <int... Indexes>
710 detail::SwizzleOp<vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
711 detail::GetOp, Indexes...>;
713 template <int... Indexes>
715 detail::SwizzleOp<const vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
716 detail::GetOp, Indexes...>;
719 template <typename... argTN>
720 using EnableIfSuitableTypes = typename detail::enable_if_t<
721 conjunction<TypeChecker<argTN, DataT>...>::value>;
723 template <typename... argTN>
724 using EnableIfSuitableNumElements = typename detail::enable_if_t<
725 SizeChecker<0, NumElements, argTN...>::value>;
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]...} {}
733 using element_type = DataT;
734 using rel_t = detail::rel_t<DataT>;
736 #ifdef __SYCL_DEVICE_ONLY__
737 using vector_t = DataType;
744 #ifdef __SYCL_DEVICE_ONLY__
745 vec(const vec &Rhs) = default;
747 constexpr vec(const vec &Rhs) : m_Data(Rhs.m_Data) {}
750 vec(vec &&Rhs) = default;
752 vec &operator=(const vec &Rhs) = default;
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,
760 operator=(const vec<rel_t, NumElements> &Rhs) {
761 *this = Rhs.template as<vec>();
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,
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,
779 template <typename Ty = DataT>
780 explicit constexpr vec(const EnableIfNotHostHalf<Ty> &arg)
781 : m_Data{(DataType)vec_data<Ty>::get(arg)} {}
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,
788 operator=(const EnableIfNotHostHalf<Ty> &Rhs) {
789 m_Data = (DataType)vec_data<Ty>::get(Rhs);
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>()} {}
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,
804 operator=(const EnableIfHostHalf<Ty> &Rhs) {
805 for (int i = 0; i < NumElements; ++i) {
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>()} {}
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,
821 operator=(const DataT &Rhs) {
822 for (int i = 0; i < NumElements; ++i) {
829 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
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,
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,
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)} {}
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>()} {}
890 #ifndef __SYCL_DEVICE_ONLY__
891 for (int I = 0; I < NumElements; ++I) {
892 std::cout << " " << I << ": " << getValue(I) << std::endl;
894 std::cout << std::endl;
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; }
908 template <int N = NumElements>
909 operator typename detail::enable_if_t<N == 1, DataT>() const {
910 return vec_data<DataT>::get(m_Data);
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); }
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) {
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)))));
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");
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");
949 detail::memcpy(&Result.m_Data, &m_Data, sizeof(decltype(Result.m_Data)));
953 template <int... SwizzleIndexes> Swizzle<SwizzleIndexes...> swizzle() {
957 template <int... SwizzleIndexes>
958 ConstSwizzle<SwizzleIndexes...> swizzle() const {
974 const DataT &operator[](int i) const {
975 return reinterpret_cast<const DataT *>(&m_Data)[i];
978 DataT &operator[](int i) { return reinterpret_cast<DataT *>(&m_Data)[i]; }
987 template <int Index> struct Indexer {
988 static constexpr int value = Index;
992 #ifdef __SYCL_ACCESS_RETURN
993 #error "Undefine __SYCL_ACCESS_RETURN macro"
995 #define __SYCL_ACCESS_RETURN this
996 #include "swizzles.def"
997 #undef __SYCL_ACCESS_RETURN
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));
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);
1012 template <int Dimensions, access::mode Mode,
1013 access::placeholder IsPlaceholder, access::target Target,
1014 typename PropertyListT>
1017 accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
1019 multi_ptr<const DataT, detail::TargetToAS<Target>::AS,
1020 access::decorated::yes>
1022 load(Offset, MultiPtr);
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 +
1032 template <int Dimensions, access::mode Mode,
1033 access::placeholder IsPlaceholder, access::target Target,
1034 typename PropertyListT>
1036 store(size_t Offset,
1037 accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
1039 multi_ptr<DataT, detail::TargetToAS<Target>::AS, access::decorated::yes>
1041 store(Offset, MultiPtr);
1045 #error "Undefine __SYCL_BINOP macro"
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 { \
1053 Ret.m_Data = m_Data BINOP Rhs.m_Data; \
1056 template <typename Ty = vec> \
1057 vec operator BINOP(const EnableIfHostHalf<Ty> &Rhs) const { \
1059 for (size_t I = 0; I < NumElements; ++I) { \
1060 Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
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), \
1070 operator BINOP(const T &Rhs) const { \
1071 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1073 vec &operator OPASSIGN(const vec &Rhs) { \
1074 *this = *this BINOP Rhs; \
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); \
1084 #define __SYCL_BINOP(BINOP, OPASSIGN) \
1085 vec operator BINOP(const vec &Rhs) const { \
1087 for (size_t I = 0; I < NumElements; ++I) { \
1088 Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
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), \
1098 operator BINOP(const T &Rhs) const { \
1099 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1101 vec &operator OPASSIGN(const vec &Rhs) { \
1102 *this = *this BINOP Rhs; \
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); \
1124 __SYCL_BINOP(>>, >>=)
1125 __SYCL_BINOP(<<, <<=)
1127 #undef __SYCL_BINOP_HELP
1136 #ifdef __SYCL_RELLOGOP
1137 #error "Undefine __SYCL_RELLOGOP macro"
1141 #ifdef __SYCL_DEVICE_ONLY__
1142 #define __SYCL_RELLOGOP(RELLOGOP) \
1143 vec<rel_t, NumElements> operator RELLOGOP(const vec &Rhs) const { \
1145 vec<rel_t, NumElements>((typename vec<rel_t, NumElements>::vector_t)( \
1146 m_Data RELLOGOP Rhs.m_Data)); \
1147 if (NumElements == 1) \
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)); \
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)))); \
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)); \
1190 #undef __SYCL_RELLOGOP
1193 #error "Undefine __SYCL_UOP macro"
1195 #define __SYCL_UOP(UOP, OPASSIGN) \
1196 vec &operator UOP() { \
1197 *this OPASSIGN vec_data<DataT>::get(1); \
1200 vec operator UOP(int) { \
1202 *this OPASSIGN vec_data<DataT>::get(1); \
1212 template <typename T = DataT>
1213 typename detail::enable_if_t<std::is_integral<vec_data_t<T>>::value, vec>
1217 #ifdef __SYCL_DEVICE_ONLY__
1218 return vec{(typename vec::DataType) ~m_Data};
1221 for (size_t I = 0; I < NumElements; ++I) {
1222 Ret.setValue(I, ~getValue(I));
1228 vec<rel_t, NumElements> operator!() const {
1231 #ifdef __SYCL_DEVICE_ONLY__
1232 return vec<rel_t, NumElements>{
1233 (typename vec<rel_t, NumElements>::DataType) !m_Data};
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)));
1243 vec operator+() const {
1246 #ifdef __SYCL_DEVICE_ONLY__
1247 return vec{+m_Data};
1250 for (size_t I = 0; I < NumElements; ++I)
1251 Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1256 vec operator-() const {
1259 #ifdef __SYCL_DEVICE_ONLY__
1260 return vec{-m_Data};
1263 for (size_t I = 0; I < NumElements; ++I)
1264 Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
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);
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)));
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)));
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,
1320 m_Data[Index] = vec_data<DataT>::get(Value);
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]);
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);
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]);
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);
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]);
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);
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);
1367 constexpr void setValue(int Index, const DataT &Value) {
1368 if (NumElements == 1)
1369 setValue(Index, Value, 0);
1371 setValue(Index, Value, 0.f);
1374 DataT getValue(int Index) const {
1375 return (NumElements == 1) ? getValue(Index, 0) : getValue(Index, 0.f);
1384 __SYCL_ALIGNED_VAR(DataType,
1385 (detail::vector_alignment<DataT, NumElements>::value),
1389 template <typename T1, typename T2, typename T3, template <typename> class T4,
1391 friend class detail::SwizzleOp;
1392 template <typename T1, int T2> friend class vec;
1395 #ifdef __cpp_deduction_guides
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>;
1406 template <typename VecT, typename OperationLeftT, typename OperationRightT,
1407 template <typename> class OperationCurrentT, int... Indexes>
1409 using DataT = typename VecT::element_type;
1411 typename std::common_type<typename OperationLeftT::DataT,
1412 typename OperationRightT::DataT>::type;
1413 static constexpr int getNumElements() { return sizeof...(Indexes); }
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)>;
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_...>;
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_...>;
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_...>;
1440 template <int IdxNum, typename T = void>
1441 using EnableIfOneIndex = typename detail::enable_if_t<
1442 1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1444 template <int IdxNum, typename T = void>
1445 using EnableIfMultipleIndexes = typename detail::enable_if_t<
1446 1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
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)>;
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)>;
1460 template <int... Indices>
1462 SwizzleOp<VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1464 template <int... Indices>
1465 using ConstSwizzle =
1466 SwizzleOp<const VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
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(); }
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>();
1480 template <int Num = getNumElements()> size_t byte_size() const noexcept {
1481 return sizeof(DataT) * (Num == 3 ? 4 : Num);
1484 template <typename T, int IdxNum = getNumElements(),
1485 typename = EnableIfOneIndex<IdxNum>,
1486 typename = EnableIfScalarType<T>>
1487 operator T() const {
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);
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);
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);
1513 #ifdef __SYCL_OPASSIGN
1514 #error "Undefine __SYCL_OPASSIGN macro."
1516 #define __SYCL_OPASSIGN(OPASSIGN, OP) \
1517 SwizzleOp &operator OPASSIGN(const DataT &Rhs) { \
1518 operatorHelper<OP>(vec_t(Rhs)); \
1521 template <typename RhsOperation> \
1522 SwizzleOp &operator OPASSIGN(const RhsOperation &Rhs) { \
1523 operatorHelper<OP>(Rhs); \
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
1540 #error "Undefine __SYCL_UOP macro"
1542 #define __SYCL_UOP(UOP, OPASSIGN) \
1543 SwizzleOp &operator UOP() { \
1544 *this OPASSIGN static_cast<DataT>(1); \
1547 vec_t operator UOP(int) { \
1548 vec_t Ret = *this; \
1549 *this OPASSIGN static_cast<DataT>(1); \
1557 template <typename T = DataT>
1558 typename detail::enable_if_t<std::is_integral<vec_data_t<T>>::value, vec_t>
1564 vec_rel_t operator!() {
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));
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);
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);
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));
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,
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));
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);
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));
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);
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));
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,
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));
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,
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));
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);
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));
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,
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));
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);
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));
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);
1730 template <typename T1, typename T2, typename T3, template <typename> class T4,
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));
1742 template <typename T1, typename T2, typename T3, template <typename> class T4,
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));
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));
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);
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));
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);
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));
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);
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));
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);
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));
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);
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));
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);
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));
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);
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));
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);
1871 template <int Index> struct Indexer {
1872 static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...};
1873 static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index];
1877 #ifdef __SYCL_ACCESS_RETURN
1878 #error "Undefine __SYCL_ACCESS_RETURN macro"
1880 #define __SYCL_ACCESS_RETURN m_Vector
1881 #include "swizzles.def"
1882 #undef __SYCL_ACCESS_RETURN
1887 template <access::address_space Space, access::decorated DecorateAddress>
1888 void load(size_t offset, multi_ptr<DataT, Space, DecorateAddress> ptr) {
1890 Tmp.template load(offset, ptr);
1894 template <typename convertT, rounding_mode roundingMode>
1895 vec<convertT, sizeof...(Indexes)> convert() const {
1898 return Tmp.template convert<convertT, roundingMode>();
1901 template <typename asT> asT as() const {
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");
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>();
1915 SwizzleOp(const SwizzleOp &Rhs)
1916 : m_Vector(Rhs.m_Vector), m_LeftOperation(Rhs.m_LeftOperation),
1917 m_RightOperation(Rhs.m_RightOperation) {}
1919 SwizzleOp(VecT *Vector, OperationLeftT LeftOperation,
1920 OperationRightT RightOperation)
1921 : m_Vector(Vector), m_LeftOperation(LeftOperation),
1922 m_RightOperation(RightOperation) {}
1924 SwizzleOp(VecT *Vector) : m_Vector(Vector) {}
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)) {}
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]);
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))));
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]);
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))));
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);
1973 OperationLeftT m_LeftOperation;
1974 OperationRightT m_RightOperation;
1977 template <typename T1, int T2> friend class sycl::vec;
1979 template <typename T1, typename T2, typename T3, template <typename> class T4,
1981 friend class SwizzleOp;
1989 #error "Undefine __SYCL_BINOP macro"
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, \
1997 operator BINOP(const T &Lhs, const vec<T, Num> &Rhs) { \
1998 return vec<T, Num>(Lhs) BINOP Rhs; \
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), \
2011 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2012 OperationCurrentT, Indexes...> &Rhs) { \
2013 vec<T1, Num> Tmp = Rhs; \
2014 return Lhs BINOP Tmp; \
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; \
2042 #ifdef __SYCL_RELLOGOP
2043 #error "Undefine __SYCL_RELLOGOP macro"
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; \
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( \
2066 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2067 OperationCurrentT, Indexes...> &Rhs) { \
2068 vec<T1, Num> Tmp = Rhs; \
2069 return Lhs RELLOGOP Tmp; \
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; \
2092 #undef __SYCL_RELLOGOP
2098 __SYCL_INLINE_VER_NAMESPACE(_V1) {
2103 constexpr bool isValidVectorSize(int N) {
2104 return N == 2 || N == 3 || N == 4 || N == 8 || N == 16;
2106 template <typename T, int N, typename V> struct VecStorage {
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");
2114 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
2115 template <typename T, int N> struct VecStorageImpl {
2116 using DataType = T __attribute__((ext_vector_type(N)));
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; \
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)
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
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>>>;
2154 template <> struct VecStorage<bool, 1, void> {
2155 using DataType = bool;
2159 struct VecStorage<bool, N, typename std::enable_if_t<isValidVectorSize(N)>> {
2161 typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
2162 std::int32_t, std::int64_t>,
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,
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>;
2178 template <typename T>
2181 typename std::enable_if_t<!is_half<T>::value && is_sgenfloat<T>::value>> {
2183 select_apply_cl_t<T, std::false_type, std::false_type, float, double>;
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))>> {
2193 typename VecStorageImpl<typename VecStorage<T, 1>::DataType, N>::DataType;
2196 template <> struct VecStorage<half, 1, void> {
2197 using DataType = sycl::detail::half_impl::StorageT;
2200 #define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
2201 template <> struct VecStorage<half, Num, void> { \
2202 using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
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
2215 #define SYCL_DEVICE_COPYABLE 1
2224 template <typename T, typename = void>
2225 struct is_device_copyable : std::false_type {};
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 {};
2238 template <typename T>
2239 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
2242 template <typename T>
2243 struct is_device_copyable<std::array<T, 0>> : std::true_type {};
2246 template <typename T, std::size_t N>
2247 struct is_device_copyable<
2249 std::enable_if_t<!std::is_trivially_copyable<std::array<T, N>>::value>>
2250 : is_device_copyable<T> {};
2253 template <typename T>
2254 struct is_device_copyable<
2256 std::enable_if_t<!std::is_trivially_copyable<std::optional<T>>::value>>
2257 : is_device_copyable<T> {};
2261 template <typename T1, typename T2>
2262 struct is_device_copyable<
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> {};
2269 template <> struct is_device_copyable<std::tuple<>> : std::true_type {};
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> {};
2281 template <> struct is_device_copyable<std::variant<>> : std::true_type {};
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 && ...)> {};
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 {};
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> {};
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>> {};
2313 template <typename T, typename = void>
2314 struct IsDeprecatedDeviceCopyable : std::false_type {};
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 {};
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> {};
2329 #ifdef __SYCL_DEVICE_ONLY__
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");
2341 template <typename T> struct CheckFieldsAreDeviceCopyable<T, 0> {};
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");
2354 template <typename T> struct CheckBasesAreDeviceCopyable<T, 0> {};
2367 template <typename FuncT>
2368 struct CheckDeviceCopyable
2369 : CheckFieldsAreDeviceCopyable<FuncT, __builtin_num_fields(FuncT)>,
2370 CheckBasesAreDeviceCopyable<FuncT, __builtin_num_bases(FuncT)> {};
2374 template <typename TransformedArgType, int Dims, typename KernelType>
2375 struct CheckDeviceCopyable<
2376 RoundedRangeKernel<TransformedArgType, Dims, KernelType>>
2377 : CheckDeviceCopyable<KernelType> {};
2379 template <typename TransformedArgType, int Dims, typename KernelType>
2380 struct CheckDeviceCopyable<
2381 RoundedRangeKernelWithKH<TransformedArgType, Dims, KernelType>>
2382 : CheckDeviceCopyable<KernelType> {};
2390 #undef __SYCL_ALIGNED_VAR