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__
61 #ifndef __SYCL_DEVICE_ONLY__
73 static constexpr
int x = 0;
74 static constexpr
int y = 1;
75 static constexpr
int z = 2;
76 static constexpr
int w = 3;
77 static constexpr
int r = 0;
78 static constexpr
int g = 1;
79 static constexpr
int b = 2;
80 static constexpr
int a = 3;
81 static constexpr
int s0 = 0;
82 static constexpr
int s1 = 1;
83 static constexpr
int s2 = 2;
84 static constexpr
int s3 = 3;
85 static constexpr
int s4 = 4;
86 static constexpr
int s5 = 5;
87 static constexpr
int s6 = 6;
88 static constexpr
int s7 = 7;
89 static constexpr
int s8 = 8;
90 static constexpr
int s9 = 9;
91 static constexpr
int sA = 10;
92 static constexpr
int sB = 11;
93 static constexpr
int sC = 12;
94 static constexpr
int sD = 13;
95 static constexpr
int sE = 14;
96 static constexpr
int sF = 15;
106 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
108 using RetType = std::uint8_t;
109 static constexpr RetType
get(
std::byte value) {
return (RetType)value; }
116 template <
typename VecT,
typename OperationLeftT,
typename OperationRightT,
117 template <
typename>
class OperationCurrentT,
int... Indexes>
123 template <
typename DataT>
157 return (Lhs == Rhs) ? -1 : 0;
163 return (Lhs != Rhs) ? -1 : 0;
169 return (Lhs >= Rhs) ? -1 : 0;
175 return (Lhs <= Rhs) ? -1 : 0;
181 return (Lhs > Rhs) ? -1 : 0;
187 return (Lhs < Rhs) ? -1 : 0;
193 return (Lhs && Rhs) ? -1 : 0;
199 return (Lhs || Rhs) ? -1 : 0;
215 template <
typename T,
typename R>
217 std::integral_constant<bool, std::is_integral<T>::value &&
218 std::is_integral<R>::value>;
220 template <
typename T,
typename R>
222 std::integral_constant<bool, is_sigeninteger<T>::value &&
225 template <
typename T,
typename R>
227 std::integral_constant<bool, is_sugeninteger<T>::value &&
230 template <
typename T,
typename R>
235 template <
typename T,
typename R>
237 std::integral_constant<bool, std::is_integral<T>::value &&
238 !(std::is_unsigned<T>::value) &&
241 template <
typename T,
typename R>
243 std::integral_constant<bool, std::is_unsigned<T>::value &&
246 template <
typename T,
typename R>
248 std::integral_constant<bool, std::is_integral<T>::value &&
251 template <
typename T,
typename R>
253 std::integral_constant<bool, detail::is_floating_point<T>::value &&
254 std::is_integral<R>::value>;
256 template <
typename T,
typename R>
258 std::integral_constant<bool, detail::is_floating_point<T>::value &&
260 template <
typename T>
262 std::integral_constant<bool, detail::is_sgentype<T>::value>;
264 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
270 #ifndef __SYCL_DEVICE_ONLY__
274 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
282 return static_cast<R
>(Value);
286 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
289 switch (roundingMode) {
291 case rounding_mode::automatic:
293 case rounding_mode::rte: {
294 int OldRoundingDirection = std::fegetround();
295 int Err = std::fesetround(FE_TONEAREST);
297 throw runtime_error(
"Unable to set rounding mode to FE_TONEAREST",
300 Err = std::fesetround(OldRoundingDirection);
306 case rounding_mode::rtz:
309 case rounding_mode::rtp:
312 case rounding_mode::rtn:
315 assert(
false &&
"Unsupported rounding mode!");
316 return static_cast<R
>(Value);
320 template <rounding_mode Mode>
322 Mode == rounding_mode::rte>;
324 template <rounding_mode Mode>
327 template <rounding_mode Mode>
330 template <rounding_mode Mode>
334 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
337 !std::is_same<T, R>::value && std::is_same<OpenCLT, OpenCLR>::value, R>
339 return static_cast<R
>(Value);
343 #define __SYCL_GENERATE_CONVERT_IMPL(DestType) \
344 template <typename T, typename R, rounding_mode roundingMode, \
345 typename OpenCLT, typename OpenCLR> \
346 detail::enable_if_t<is_sint_to_sint<T, R>::value && \
347 !std::is_same<OpenCLT, OpenCLR>::value && \
348 (std::is_same<OpenCLR, cl_##DestType>::value || \
349 (std::is_same<OpenCLR, signed char>::value && \
350 std::is_same<DestType, char>::value)), \
352 convertImpl(T Value) { \
353 OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
354 return __spirv_SConvert##_R##DestType(OpValue); \
357 __SYCL_GENERATE_CONVERT_IMPL(
char)
358 __SYCL_GENERATE_CONVERT_IMPL(
short)
359 __SYCL_GENERATE_CONVERT_IMPL(
int)
360 __SYCL_GENERATE_CONVERT_IMPL(
long)
362 #undef __SYCL_GENERATE_CONVERT_IMPL
365 #define __SYCL_GENERATE_CONVERT_IMPL(DestType) \
366 template <typename T, typename R, rounding_mode roundingMode, \
367 typename OpenCLT, typename OpenCLR> \
368 detail::enable_if_t<is_uint_to_uint<T, R>::value && \
369 !std::is_same<OpenCLT, OpenCLR>::value && \
370 std::is_same<OpenCLR, cl_##DestType>::value, \
372 convertImpl(T Value) { \
373 OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
374 return __spirv_UConvert##_R##DestType(OpValue); \
377 __SYCL_GENERATE_CONVERT_IMPL(
uchar)
378 __SYCL_GENERATE_CONVERT_IMPL(
ushort)
379 __SYCL_GENERATE_CONVERT_IMPL(
uint)
380 __SYCL_GENERATE_CONVERT_IMPL(
ulong)
382 #undef __SYCL_GENERATE_CONVERT_IMPL
385 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
387 detail::enable_if_t<is_sint_to_from_uint<T, R>::value &&
388 is_standard_type<OpenCLT>::value &&
389 is_standard_type<OpenCLR>::value,
392 return static_cast<R
>(Value);
396 #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
397 template <typename T, typename R, rounding_mode roundingMode, \
398 typename OpenCLT, typename OpenCLR> \
399 detail::enable_if_t<is_sint_to_float<T, R>::value && \
400 (std::is_same<OpenCLR, DestType>::value || \
401 (std::is_same<OpenCLR, _Float16>::value && \
402 std::is_same<DestType, half>::value)), \
404 convertImpl(T Value) { \
405 OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
406 return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
409 __SYCL_GENERATE_CONVERT_IMPL(SToF,
half)
410 __SYCL_GENERATE_CONVERT_IMPL(SToF,
float)
411 __SYCL_GENERATE_CONVERT_IMPL(SToF,
double)
413 #undef __SYCL_GENERATE_CONVERT_IMPL
416 #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType) \
417 template <typename T, typename R, rounding_mode roundingMode, \
418 typename OpenCLT, typename OpenCLR> \
419 detail::enable_if_t<is_uint_to_float<T, R>::value && \
420 (std::is_same<OpenCLR, DestType>::value || \
421 (std::is_same<OpenCLR, _Float16>::value && \
422 std::is_same<DestType, half>::value)), \
424 convertImpl(T Value) { \
425 OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
426 return __spirv_Convert##SPIRVOp##_R##DestType(OpValue); \
429 __SYCL_GENERATE_CONVERT_IMPL(UToF,
half)
430 __SYCL_GENERATE_CONVERT_IMPL(UToF,
float)
431 __SYCL_GENERATE_CONVERT_IMPL(UToF,
double)
433 #undef __SYCL_GENERATE_CONVERT_IMPL
436 #define __SYCL_GENERATE_CONVERT_IMPL(DestType, RoundingMode, \
437 RoundingModeCondition) \
438 template <typename T, typename R, rounding_mode roundingMode, \
439 typename OpenCLT, typename OpenCLR> \
440 detail::enable_if_t<is_float_to_float<T, R>::value && \
441 !std::is_same<OpenCLT, OpenCLR>::value && \
442 (std::is_same<OpenCLR, DestType>::value || \
443 (std::is_same<OpenCLR, _Float16>::value && \
444 std::is_same<DestType, half>::value)) && \
445 RoundingModeCondition<roundingMode>::value, \
447 convertImpl(T Value) { \
448 OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
449 return __spirv_FConvert##_R##DestType##_##RoundingMode(OpValue); \
452 #define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \
453 RoundingModeCondition) \
454 __SYCL_GENERATE_CONVERT_IMPL(double, RoundingMode, RoundingModeCondition) \
455 __SYCL_GENERATE_CONVERT_IMPL(float, RoundingMode, RoundingModeCondition) \
456 __SYCL_GENERATE_CONVERT_IMPL(half, RoundingMode, RoundingModeCondition)
458 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic)
459 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz)
460 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp)
461 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)
463 #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
464 #undef __SYCL_GENERATE_CONVERT_IMPL
467 #define __SYCL_GENERATE_CONVERT_IMPL(SPIRVOp, DestType, RoundingMode, \
468 RoundingModeCondition) \
469 template <typename T, typename R, rounding_mode roundingMode, \
470 typename OpenCLT, typename OpenCLR> \
471 detail::enable_if_t<is_float_to_int<T, R>::value && \
472 (std::is_same<OpenCLR, cl_##DestType>::value || \
473 (std::is_same<OpenCLR, signed char>::value && \
474 std::is_same<DestType, char>::value)) && \
475 RoundingModeCondition<roundingMode>::value, \
477 convertImpl(T Value) { \
478 OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value); \
479 return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue); \
482 #define __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(RoundingMode, \
483 RoundingModeCondition) \
484 __SYCL_GENERATE_CONVERT_IMPL(FToS, int, RoundingMode, RoundingModeCondition) \
485 __SYCL_GENERATE_CONVERT_IMPL(FToS, char, RoundingMode, \
486 RoundingModeCondition) \
487 __SYCL_GENERATE_CONVERT_IMPL(FToS, short, RoundingMode, \
488 RoundingModeCondition) \
489 __SYCL_GENERATE_CONVERT_IMPL(FToS, long, RoundingMode, \
490 RoundingModeCondition) \
491 __SYCL_GENERATE_CONVERT_IMPL(FToU, uint, RoundingMode, \
492 RoundingModeCondition) \
493 __SYCL_GENERATE_CONVERT_IMPL(FToU, uchar, RoundingMode, \
494 RoundingModeCondition) \
495 __SYCL_GENERATE_CONVERT_IMPL(FToU, ushort, RoundingMode, \
496 RoundingModeCondition) \
497 __SYCL_GENERATE_CONVERT_IMPL(FToU, ulong, RoundingMode, RoundingModeCondition)
499 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rte, RteOrAutomatic)
500 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtz, Rtz)
501 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtp, Rtp)
502 __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)
504 #undef __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE
505 #undef __SYCL_GENERATE_CONVERT_IMPL
508 template <
typename T,
typename R,
rounding_mode roundingMode,
typename OpenCLT,
511 ((!is_standard_type<T>::value && !is_standard_type<OpenCLT>::value) ||
512 (!is_standard_type<R>::value && !is_standard_type<OpenCLR>::value)) &&
513 !std::is_same<OpenCLT, OpenCLR>::value,
516 return static_cast<R
>(Value);
519 #endif // __SYCL_DEVICE_ONLY__
522 template <
typename TransformedArgType,
int Dims,
typename KernelType>
523 class RoundedRangeKernel;
524 template <
typename TransformedArgType,
int Dims,
typename KernelType>
525 class RoundedRangeKernelWithKH;
531 template <
typename T>
534 #if defined(_WIN32) && (_MSC_VER)
545 #pragma message("Alignment of class vec is not in accordance with SYCL \
546 specification requirements, a limitation of the MSVC compiler(Error C2719).\
547 Requested alignment applied, limited at 64.")
548 #define __SYCL_ALIGNED_VAR(type, x, var) \
549 type __declspec(align((x < 64) ? x : 64)) var
551 #define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
558 template <
typename Type,
int NumElements>
class vec {
564 typename detail::BaseCLTypeConverter<DataT, NumElements>::DataType;
566 static constexpr
int getNumElements() {
return NumElements; }
569 template <
int Counter,
int MaxValue,
class...>
571 std::true_type, std::false_type> {};
573 template <
int Counter,
int MaxValue,
typename DataT_,
class... tail>
574 struct SizeChecker<Counter, MaxValue, DataT_, tail...>
576 SizeChecker<Counter + 1, MaxValue, tail...>,
579 #define __SYCL_ALLOW_VECTOR_SIZES(num_elements) \
580 template <int Counter, int MaxValue, typename DataT_, class... tail> \
581 struct SizeChecker<Counter, MaxValue, vec<DataT_, num_elements>, tail...> \
582 : detail::conditional_t< \
583 Counter + (num_elements) <= MaxValue, \
584 SizeChecker<Counter + (num_elements), MaxValue, tail...>, \
585 std::false_type> {}; \
586 template <int Counter, int MaxValue, typename DataT_, typename T2, \
587 typename T3, template <typename> class T4, int... T5, \
589 struct SizeChecker< \
591 detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
593 : detail::conditional_t< \
594 Counter + sizeof...(T5) <= MaxValue, \
595 SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
596 std::false_type> {}; \
597 template <int Counter, int MaxValue, typename DataT_, typename T2, \
598 typename T3, template <typename> class T4, int... T5, \
600 struct SizeChecker< \
602 detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
604 : detail::conditional_t< \
605 Counter + sizeof...(T5) <= MaxValue, \
606 SizeChecker<Counter + sizeof...(T5), MaxValue, tail...>, \
609 __SYCL_ALLOW_VECTOR_SIZES(1)
610 __SYCL_ALLOW_VECTOR_SIZES(2)
611 __SYCL_ALLOW_VECTOR_SIZES(3)
612 __SYCL_ALLOW_VECTOR_SIZES(4)
613 __SYCL_ALLOW_VECTOR_SIZES(8)
614 __SYCL_ALLOW_VECTOR_SIZES(16)
615 #undef __SYCL_ALLOW_VECTOR_SIZES
617 template <class...> struct conjunction : std::true_type {};
618 template <class B1, class... tail>
619 struct conjunction<B1, tail...>
620 : detail::conditional_t<bool(B1::value), conjunction<tail...>, B1> {};
623 template <typename T, typename DataT_>
624 struct TypeChecker : std::is_convertible<T, DataT_> {};
625 #define __SYCL_ALLOW_VECTOR_TYPES(num_elements) \
626 template <typename DataT_> \
627 struct TypeChecker<vec<DataT_, num_elements>, DataT_> : std::true_type {}; \
628 template <typename DataT_, typename T2, typename T3, \
629 template <typename> class T4, int... T5> \
630 struct TypeChecker< \
631 detail::SwizzleOp<vec<DataT_, num_elements>, T2, T3, T4, T5...>, DataT_> \
632 : std::true_type {}; \
633 template <typename DataT_, typename T2, typename T3, \
634 template <typename> class T4, int... T5> \
635 struct TypeChecker< \
636 detail::SwizzleOp<const vec<DataT_, num_elements>, T2, T3, T4, T5...>, \
637 DataT_> : std::true_type {};
639 __SYCL_ALLOW_VECTOR_TYPES(1)
640 __SYCL_ALLOW_VECTOR_TYPES(2)
641 __SYCL_ALLOW_VECTOR_TYPES(3)
642 __SYCL_ALLOW_VECTOR_TYPES(4)
643 __SYCL_ALLOW_VECTOR_TYPES(8)
644 __SYCL_ALLOW_VECTOR_TYPES(16)
645 #undef __SYCL_ALLOW_VECTOR_TYPES
647 template <int... Indexes>
649 detail::SwizzleOp<vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
650 detail::GetOp, Indexes...>;
652 template <int... Indexes>
654 detail::SwizzleOp<const vec, detail::GetOp<DataT>, detail::GetOp<DataT>,
655 detail::GetOp, Indexes...>;
658 template <typename... argTN>
659 using EnableIfSuitableTypes = typename detail::enable_if_t<
660 conjunction<TypeChecker<argTN, DataT>...>::value>;
662 template <typename... argTN>
663 using EnableIfSuitableNumElements = typename detail::enable_if_t<
664 SizeChecker<0, NumElements, argTN...>::value>;
667 using element_type = DataT;
668 using rel_t = detail::rel_t<DataT>;
670 #ifdef __SYCL_DEVICE_ONLY__
671 using vector_t = DataType;
678 #ifdef __SYCL_DEVICE_ONLY__
679 vec(const vec &Rhs) = default;
681 constexpr vec(const vec &Rhs) : m_Data(Rhs.m_Data) {}
684 vec(vec &&Rhs) = default;
686 vec &operator=(const vec &Rhs) = default;
689 template <typename Ty = DataT>
690 typename detail::enable_if_t<
691 !std::is_same<Ty, rel_t>::value &&
692 std::is_convertible<vec_data_t<Ty>, rel_t>::value,
694 operator=(const vec<rel_t, NumElements> &Rhs) {
695 *this = Rhs.template as<vec>();
699 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
700 template <typename T = void>
701 using EnableIfNotHostHalf = typename detail::enable_if_t<
702 !std::is_same<DataT, cl::sycl::detail::half_impl::half>::value ||
703 !std::is_same<cl::sycl::detail::half_impl::StorageT,
704 cl::sycl::detail::host_half_impl::half_v2>::value,
706 template <typename T = void>
707 using EnableIfHostHalf = typename detail::enable_if_t<
708 std::is_same<DataT, cl::sycl::detail::half_impl::half>::value &&
709 std::is_same<cl::sycl::detail::half_impl::StorageT,
710 cl::sycl::detail::host_half_impl::half_v2>::value,
713 template <typename Ty = DataT>
714 explicit constexpr vec(const EnableIfNotHostHalf<Ty> &arg) {
715 m_Data = (DataType)vec_data<Ty>::get(arg);
718 template <typename Ty = DataT>
719 typename detail::enable_if_t<
720 std::is_fundamental<vec_data_t<Ty>>::value ||
721 std::is_same<typename detail::remove_const_t<Ty>, half>::value,
723 operator=(const EnableIfNotHostHalf<Ty> &Rhs) {
724 m_Data = (DataType)vec_data<Ty>::get(Rhs);
728 template <typename Ty = DataT>
729 explicit constexpr vec(const EnableIfHostHalf<Ty> &arg) {
730 for (int i = 0; i < NumElements; ++i) {
735 template <typename Ty = DataT>
736 typename detail::enable_if_t<
737 std::is_fundamental<vec_data_t<Ty>>::value ||
738 std::is_same<typename detail::remove_const_t<Ty>, half>::value,
740 operator=(const EnableIfHostHalf<Ty> &Rhs) {
741 for (int i = 0; i < NumElements; ++i) {
747 explicit constexpr vec(const DataT &arg) {
748 for (int i = 0; i < NumElements; ++i) {
753 template <typename Ty = DataT>
754 typename detail::enable_if_t<
755 std::is_fundamental<vec_data_t<Ty>>::value ||
756 std::is_same<typename detail::remove_const_t<Ty>, half>::value,
758 operator=(const DataT &Rhs) {
759 for (int i = 0; i < NumElements; ++i) {
766 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
772 template <int IdxNum, typename T = void>
773 using EnableIfMultipleElems = typename detail::enable_if_t<
774 std::is_convertible<T, DataT>::value && NumElements == IdxNum, DataT>;
775 template <typename Ty = DataT>
776 constexpr vec(const EnableIfMultipleElems<2, Ty> Arg0,
777 const EnableIfNotHostHalf<Ty> Arg1)
778 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1)} {}
779 template <typename Ty = DataT>
780 constexpr vec(const EnableIfMultipleElems<3, Ty> Arg0,
781 const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2)
782 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
783 vec_data<Ty>::get(Arg2)} {}
784 template <typename Ty = DataT>
785 constexpr vec(const EnableIfMultipleElems<4, Ty> Arg0,
786 const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2,
788 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
789 vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3)} {}
790 template <typename Ty = DataT>
791 constexpr vec(const EnableIfMultipleElems<8, Ty> Arg0,
792 const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2,
793 const DataT Arg3, const DataT Arg4, const DataT Arg5,
794 const DataT Arg6, const DataT Arg7)
795 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
796 vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3),
797 vec_data<Ty>::get(Arg4), vec_data<Ty>::get(Arg5),
798 vec_data<Ty>::get(Arg6), vec_data<Ty>::get(Arg7)} {}
799 template <typename Ty = DataT>
800 constexpr vec(const EnableIfMultipleElems<16, Ty> Arg0,
801 const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2,
802 const DataT Arg3, const DataT Arg4, const DataT Arg5,
803 const DataT Arg6, const DataT Arg7, const DataT Arg8,
804 const DataT Arg9, const DataT ArgA, const DataT ArgB,
805 const DataT ArgC, const DataT ArgD, const DataT ArgE,
807 : m_Data{vec_data<Ty>::get(Arg0), vec_data<Ty>::get(Arg1),
808 vec_data<Ty>::get(Arg2), vec_data<Ty>::get(Arg3),
809 vec_data<Ty>::get(Arg4), vec_data<Ty>::get(Arg5),
810 vec_data<Ty>::get(Arg6), vec_data<Ty>::get(Arg7),
811 vec_data<Ty>::get(Arg8), vec_data<Ty>::get(Arg9),
812 vec_data<Ty>::get(ArgA), vec_data<Ty>::get(ArgB),
813 vec_data<Ty>::get(ArgC), vec_data<Ty>::get(ArgD),
814 vec_data<Ty>::get(ArgE), vec_data<Ty>::get(ArgF)} {}
819 template <typename... argTN, typename = EnableIfSuitableTypes<argTN...>,
820 typename = EnableIfSuitableNumElements<argTN...>>
821 constexpr vec(const argTN &... args) {
822 vaargCtorHelper(0, args...);
827 #ifndef __SYCL_DEVICE_ONLY__
828 for (int I = 0; I < NumElements; ++I) {
829 std::cout << " " << I << ": " << getValue(I) << std::endl;
831 std::cout << std::endl;
835 #ifdef __SYCL_DEVICE_ONLY__
836 template <typename vector_t_ = vector_t,
837 typename = typename detail::enable_if_t<
838 std::is_same<vector_t_, vector_t>::value &&
839 !std::is_same<vector_t_, DataT>::value>>
840 constexpr vec(vector_t openclVector) : m_Data(openclVector) {}
841 operator vector_t() const { return m_Data; }
845 template <int N = NumElements>
846 operator typename detail::enable_if_t<N == 1, DataT>() const {
847 return vec_data<DataT>::get(m_Data);
850 __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
851 static constexpr size_t get_count() { return size(); }
852 static constexpr size_t size() noexcept { return NumElements; }
853 __SYCL2020_DEPRECATED(
854 "get_size() is deprecated, please use byte_size() instead")
855 static constexpr size_t get_size() { return byte_size(); }
856 static constexpr size_t byte_size() { return sizeof(m_Data); }
858 template <typename convertT,
859 rounding_mode roundingMode = rounding_mode::automatic>
860 vec<convertT, NumElements> convert() const {
861 static_assert(std::is_integral<vec_data_t<convertT>>::value ||
862 detail::is_floating_point<convertT>::value,
863 "Unsupported convertT");
864 vec<convertT, NumElements> Result;
865 using OpenCLT = detail::ConvertToOpenCLType_t<vec_data_t<DataT>>;
866 using OpenCLR = detail::ConvertToOpenCLType_t<vec_data_t<convertT>>;
867 for (size_t I = 0; I < NumElements; ++I) {
869 I, vec_data<convertT>::get(
870 detail::convertImpl<vec_data_t<DataT>, vec_data_t<convertT>,
871 roundingMode, OpenCLT, OpenCLR>(
872 vec_data<DataT>::get(getValue(I)))));
877 template <typename asT> asT as() const {
878 static_assert((sizeof(*this) == sizeof(asT)),
879 "The new SYCL vec type must have the same storage size in "
880 "bytes as this SYCL vec");
882 detail::is_contained<asT, detail::gtl::vector_basic_list>::value,
883 "asT must be SYCL vec of a different element type and "
884 "number of elements specified by asT");
886 detail::memcpy(&Result.m_Data, &m_Data, sizeof(decltype(Result.m_Data)));
890 template <int... SwizzleIndexes> Swizzle<SwizzleIndexes...> swizzle() {
894 template <int... SwizzleIndexes>
895 ConstSwizzle<SwizzleIndexes...> swizzle() const {
911 const DataT &operator[](int i) const {
912 return reinterpret_cast<const DataT *>(&m_Data)[i];
915 DataT &operator[](int i) { return reinterpret_cast<DataT *>(&m_Data)[i]; }
926 static constexpr int value = Index;
930 #ifdef __SYCL_ACCESS_RETURN
931 #error "Undefine __SYCL_ACCESS_RETURN macro"
933 #define __SYCL_ACCESS_RETURN this
934 #include "swizzles.def"
935 #undef __SYCL_ACCESS_RETURN
938 template <access::address_space Space>
939 void load(size_t Offset, multi_ptr<const DataT, Space> Ptr) {
940 for (int I = 0; I < NumElements; I++) {
942 *multi_ptr<const DataT, Space>(Ptr + Offset * NumElements + I));
945 template <access::address_space Space>
946 void load(size_t Offset, multi_ptr<DataT, Space> Ptr) {
947 multi_ptr<const DataT, Space> ConstPtr(Ptr);
948 load(Offset, ConstPtr);
950 template <int Dimensions, access::mode Mode,
951 access::placeholder IsPlaceholder, access::target Target,
952 typename PropertyListT>
955 accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
957 multi_ptr<const DataT, detail::TargetToAS<Target>::AS> MultiPtr(Acc);
958 load(Offset, MultiPtr);
960 template <access::address_space Space>
961 void store(size_t Offset, multi_ptr<DataT, Space> Ptr) const {
962 for (int I = 0; I < NumElements; I++) {
963 *multi_ptr<DataT, Space>(Ptr + Offset * NumElements + I) = getValue(I);
966 template <int Dimensions, access::mode Mode,
967 access::placeholder IsPlaceholder, access::target Target,
968 typename PropertyListT>
971 accessor<DataT, Dimensions, Mode, Target, IsPlaceholder, PropertyListT>
973 multi_ptr<DataT, detail::TargetToAS<Target>::AS> MultiPtr(Acc);
974 store(Offset, MultiPtr);
978 #error "Undefine __SYCL_BINOP macro"
981 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
982 #define __SYCL_BINOP(BINOP, OPASSIGN) \
983 template <typename Ty = vec> \
984 vec operator BINOP(const EnableIfNotHostHalf<Ty> &Rhs) const { \
986 Ret.m_Data = m_Data BINOP Rhs.m_Data; \
989 template <typename Ty = vec> \
990 vec operator BINOP(const EnableIfHostHalf<Ty> &Rhs) const { \
992 for (size_t I = 0; I < NumElements; ++I) { \
993 Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
997 template <typename T> \
998 typename detail::enable_if_t< \
999 std::is_convertible<DataT, T>::value && \
1000 (std::is_fundamental<vec_data_t<T>>::value || \
1001 std::is_same<typename detail::remove_const_t<T>, half>::value), \
1003 operator BINOP(const T &Rhs) const { \
1004 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1006 vec &operator OPASSIGN(const vec &Rhs) { \
1007 *this = *this BINOP Rhs; \
1010 template <int Num = NumElements> \
1011 typename detail::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1012 const DataT &Rhs) { \
1013 *this = *this BINOP vec(Rhs); \
1017 #define __SYCL_BINOP(BINOP, OPASSIGN) \
1018 vec operator BINOP(const vec &Rhs) const { \
1020 for (size_t I = 0; I < NumElements; ++I) { \
1021 Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I))); \
1025 template <typename T> \
1026 typename detail::enable_if_t< \
1027 std::is_convertible<DataT, T>::value && \
1028 (std::is_fundamental<vec_data_t<T>>::value || \
1029 std::is_same<typename detail::remove_const_t<T>, half>::value), \
1031 operator BINOP(const T &Rhs) const { \
1032 return *this BINOP vec(static_cast<const DataT &>(Rhs)); \
1034 vec &operator OPASSIGN(const vec &Rhs) { \
1035 *this = *this BINOP Rhs; \
1038 template <int Num = NumElements> \
1039 typename detail::enable_if_t<Num != 1, vec &> operator OPASSIGN( \
1040 const DataT &Rhs) { \
1041 *this = *this BINOP vec(Rhs); \
1057 __SYCL_BINOP(>>, >>=)
1058 __SYCL_BINOP(<<, <<=)
1060 #undef __SYCL_BINOP_HELP
1069 #ifdef __SYCL_RELLOGOP
1070 #error "Undefine __SYCL_RELLOGOP macro"
1074 #ifdef __SYCL_DEVICE_ONLY__
1075 #define __SYCL_RELLOGOP(RELLOGOP) \
1076 vec<rel_t, NumElements> operator RELLOGOP(const vec &Rhs) const { \
1078 vec<rel_t, NumElements>((typename vec<rel_t, NumElements>::vector_t)( \
1079 m_Data RELLOGOP Rhs.m_Data)); \
1080 if (NumElements == 1) \
1084 template <typename T> \
1085 typename detail::enable_if_t< \
1086 std::is_convertible<T, DataT>::value && \
1087 (std::is_fundamental<vec_data_t<T>>::value || \
1088 std::is_same<T, half>::value), \
1089 vec<rel_t, NumElements>> \
1090 operator RELLOGOP(const T &Rhs) const { \
1091 return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1094 #define __SYCL_RELLOGOP(RELLOGOP) \
1095 vec<rel_t, NumElements> operator RELLOGOP(const vec &Rhs) const { \
1096 vec<rel_t, NumElements> Ret; \
1097 for (size_t I = 0; I < NumElements; ++I) { \
1098 Ret.setValue(I, -(vec_data<DataT>::get(getValue(I)) \
1099 RELLOGOP vec_data<DataT>::get(Rhs.getValue(I)))); \
1103 template <typename T> \
1104 typename detail::enable_if_t< \
1105 std::is_convertible<T, DataT>::value && \
1106 (std::is_fundamental<vec_data_t<T>>::value || \
1107 std::is_same<T, half>::value), \
1108 vec<rel_t, NumElements>> \
1109 operator RELLOGOP(const T &Rhs) const { \
1110 return *this RELLOGOP vec(static_cast<const DataT &>(Rhs)); \
1123 #undef __SYCL_RELLOGOP
1126 #error "Undefine __SYCL_UOP macro"
1128 #define __SYCL_UOP(UOP, OPASSIGN) \
1129 vec &operator UOP() { \
1130 *this OPASSIGN vec_data<DataT>::get(1); \
1133 vec operator UOP(int) { \
1135 *this OPASSIGN vec_data<DataT>::get(1); \
1145 template <typename T = DataT>
1146 typename detail::enable_if_t<std::is_integral<vec_data_t<T>>::value, vec>
1150 #ifdef __SYCL_DEVICE_ONLY__
1152 (typename vec::DataType)~m_Data};
1155 for (size_t I = 0; I < NumElements; ++I) {
1156 Ret.setValue(I, ~getValue(I));
1162 vec<rel_t, NumElements> operator!() const {
1165 #ifdef __SYCL_DEVICE_ONLY__
1166 return vec<rel_t, NumElements>{
1167 (typename vec<rel_t, NumElements>::DataType)!m_Data};
1169 vec<rel_t, NumElements> Ret;
1170 for (size_t I = 0; I < NumElements; ++I) {
1171 Ret.setValue(I, !vec_data<DataT>::get(getValue(I)));
1177 vec operator+() const {
1180 #ifdef __SYCL_DEVICE_ONLY__
1181 return vec{+m_Data};
1184 for (size_t I = 0; I < NumElements; ++I)
1185 Ret.setValue(I, vec_data<DataT>::get(+vec_data<DataT>::get(getValue(I))));
1190 vec operator-() const {
1193 #ifdef __SYCL_DEVICE_ONLY__
1194 return vec{-m_Data};
1197 for (size_t I = 0; I < NumElements; ++I)
1198 Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
1212 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1213 template <template <typename> class Operation,
1214 typename Ty = vec<DataT, NumElements>>
1215 vec<DataT, NumElements>
1216 operatorHelper(const EnableIfNotHostHalf<Ty> &Rhs) const {
1217 vec<DataT, NumElements> Result;
1218 Operation<DataType> Op;
1219 Result.m_Data = Op(m_Data, Rhs.m_Data);
1223 template <template <typename> class Operation,
1224 typename Ty = vec<DataT, NumElements>>
1225 vec<DataT, NumElements>
1226 operatorHelper(const EnableIfHostHalf<Ty> &Rhs) const {
1227 vec<DataT, NumElements> Result;
1228 Operation<DataT> Op;
1229 for (size_t I = 0; I < NumElements; ++I) {
1230 Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1235 template <template <typename> class Operation>
1236 vec<DataT, NumElements>
1237 operatorHelper(const vec<DataT, NumElements> &Rhs) const {
1238 vec<DataT, NumElements> Result;
1239 Operation<DataT> Op;
1240 for (size_t I = 0; I < NumElements; ++I) {
1241 Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
1249 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
1250 template <int Num = NumElements, typename Ty = int,
1251 typename = typename detail::enable_if_t<1 != Num>>
1252 constexpr void setValue(EnableIfNotHostHalf<Ty> Index, const DataT &Value,
1254 m_Data[Index] = vec_data<DataT>::get(Value);
1257 template <int Num = NumElements, typename Ty = int,
1258 typename = typename detail::enable_if_t<1 != Num>>
1259 DataT getValue(EnableIfNotHostHalf<Ty> Index, int) const {
1260 return vec_data<DataT>::get(m_Data[Index]);
1263 template <int Num = NumElements, typename Ty = int,
1264 typename = typename detail::enable_if_t<1 != Num>>
1265 constexpr void setValue(EnableIfHostHalf<Ty> Index, const DataT &Value, int) {
1266 m_Data.s[Index] = vec_data<DataT>::get(Value);
1269 template <int Num = NumElements, typename Ty = int,
1270 typename = typename detail::enable_if_t<1 != Num>>
1271 DataT getValue(EnableIfHostHalf<Ty> Index, int) const {
1272 return vec_data<DataT>::get(m_Data.s[Index]);
1275 template <int Num = NumElements,
1276 typename = typename detail::enable_if_t<1 != Num>>
1277 constexpr void setValue(int Index, const DataT &Value, int) {
1278 m_Data.s[Index] = vec_data<DataT>::get(Value);
1281 template <int Num = NumElements,
1282 typename = typename detail::enable_if_t<1 != Num>>
1283 DataT getValue(int Index, int) const {
1284 return vec_data<DataT>::get(m_Data.s[Index]);
1288 template <int Num = NumElements,
1289 typename = typename detail::enable_if_t<1 == Num>>
1290 constexpr void setValue(int, const DataT &Value, float) {
1291 m_Data = vec_data<DataT>::get(Value);
1294 template <int Num = NumElements,
1295 typename = typename detail::enable_if_t<1 == Num>>
1296 DataT getValue(int, float) const {
1297 return vec_data<DataT>::get(m_Data);
1301 constexpr void setValue(int Index, const DataT &Value) {
1302 if (NumElements == 1)
1303 setValue(Index, Value, 0);
1305 setValue(Index, Value, 0.f);
1308 DataT getValue(int Index) const {
1309 return (NumElements == 1) ? getValue(Index, 0) : getValue(Index, 0.f);
1313 template <typename T, typename... argTN>
1314 constexpr int vaargCtorHelper(int Idx, const T &arg) {
1319 template <typename DataT_, int NumElements_>
1320 constexpr int vaargCtorHelper(int Idx, const vec<DataT_, NumElements_> &arg) {
1321 for (size_t I = 0; I < NumElements_; ++I) {
1322 setValue(Idx + I, arg.getValue(I));
1324 return Idx + NumElements_;
1327 template <typename DataT_, int NumElements_, typename T2, typename T3,
1328 template <typename> class T4, int... T5>
1330 vaargCtorHelper(int Idx, const detail::SwizzleOp<vec<DataT_, NumElements_>,
1331 T2, T3, T4, T5...> &arg) {
1332 size_t NumElems = sizeof...(T5);
1333 for (size_t I = 0; I < NumElems; ++I) {
1334 setValue(Idx + I, arg.getValue(I));
1336 return Idx + NumElems;
1339 template <typename DataT_, int NumElements_, typename T2, typename T3,
1340 template <typename> class T4, int... T5>
1342 vaargCtorHelper(int Idx,
1343 const detail::SwizzleOp<const vec<DataT_, NumElements_>, T2,
1344 T3, T4, T5...> &arg) {
1345 size_t NumElems = sizeof...(T5);
1346 for (size_t I = 0; I < NumElems; ++I) {
1347 setValue(Idx + I, arg.getValue(I));
1349 return Idx + NumElems;
1352 template <typename T1, typename... argTN>
1353 constexpr void vaargCtorHelper(int Idx, const T1 &arg,
1354 const argTN &... args) {
1355 int NewIdx = vaargCtorHelper(Idx, arg);
1356 vaargCtorHelper(NewIdx, args...);
1359 template <typename DataT_, int NumElements_, typename... argTN>
1360 constexpr void vaargCtorHelper(int Idx, const vec<DataT_, NumElements_> &arg,
1361 const argTN &... args) {
1362 int NewIdx = vaargCtorHelper(Idx, arg);
1363 vaargCtorHelper(NewIdx, args...);
1372 __SYCL_ALIGNED_VAR(DataType,
1373 (detail::vector_alignment<DataT, NumElements>::value),
1377 template <typename T1, typename T2, typename T3, template <typename> class T4,
1379 friend class detail::SwizzleOp;
1380 template <typename T1, int T2> friend class vec;
1383 #ifdef __cpp_deduction_guides
1385 template <class T, class... U,
1386 class = detail::enable_if_t<(std::is_same<T, U>::value && ...)>>
1387 vec(T, U...)->vec<T, sizeof...(U) + 1>;
1394 template <typename VecT, typename OperationLeftT, typename OperationRightT,
1395 template <typename> class OperationCurrentT, int... Indexes>
1397 using DataT = typename VecT::element_type;
1399 typename std::common_type<typename OperationLeftT::DataT,
1400 typename OperationRightT::DataT>::type;
1401 static constexpr int getNumElements() { return sizeof...(Indexes); }
1403 using rel_t = detail::rel_t<DataT>;
1404 using vec_t = vec<DataT, sizeof...(Indexes)>;
1405 using vec_rel_t = vec<rel_t, sizeof...(Indexes)>;
1407 template <typename OperationRightT_,
1408 template <typename> class OperationCurrentT_, int... Idx_>
1409 using NewLHOp = SwizzleOp<VecT,
1410 SwizzleOp<VecT, OperationLeftT, OperationRightT,
1411 OperationCurrentT, Indexes...>,
1412 OperationRightT_, OperationCurrentT_, Idx_...>;
1414 template <typename OperationRightT_,
1415 template <typename> class OperationCurrentT_, int... Idx_>
1416 using NewRelOp = SwizzleOp<vec<rel_t, VecT::getNumElements()>,
1417 SwizzleOp<VecT, OperationLeftT, OperationRightT,
1418 OperationCurrentT, Indexes...>,
1419 OperationRightT_, OperationCurrentT_, Idx_...>;
1421 template <typename OperationLeftT_,
1422 template <typename> class OperationCurrentT_, int... Idx_>
1423 using NewRHOp = SwizzleOp<VecT, OperationLeftT_,
1424 SwizzleOp<VecT, OperationLeftT, OperationRightT,
1425 OperationCurrentT, Indexes...>,
1426 OperationCurrentT_, Idx_...>;
1428 template <int IdxNum, typename T = void>
1429 using EnableIfOneIndex = typename detail::enable_if_t<
1430 1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1432 template <int IdxNum, typename T = void>
1433 using EnableIfMultipleIndexes = typename detail::enable_if_t<
1434 1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>;
1436 template <typename T>
1437 using EnableIfScalarType = typename detail::enable_if_t<
1438 std::is_convertible<DataT, T>::value &&
1439 (std::is_fundamental<vec_data_t<T>>::value ||
1440 std::is_same<typename detail::remove_const_t<T>, half>::value)>;
1442 template <typename T>
1443 using EnableIfNoScalarType = typename detail::enable_if_t<
1444 !std::is_convertible<DataT, T>::value ||
1445 !(std::is_fundamental<vec_data_t<T>>::value ||
1446 std::is_same<typename detail::remove_const_t<T>, half>::value)>;
1448 template <int... Indices>
1450 SwizzleOp<VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1452 template <int... Indices>
1453 using ConstSwizzle =
1454 SwizzleOp<const VecT, GetOp<DataT>, GetOp<DataT>, GetOp, Indices...>;
1457 __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
1458 size_t get_count() const { return size(); }
1459 size_t size() const noexcept { return getNumElements(); }
1461 template <int Num = getNumElements()>
1462 __SYCL2020_DEPRECATED(
1463 "get_size() is deprecated, please use byte_size() instead")
1464 size_t get_size() const {
1465 return byte_size<Num>();
1468 template <int Num = getNumElements()> size_t byte_size() const noexcept {
1469 return sizeof(DataT) * (Num == 3 ? 4 : Num);
1472 template <typename T, int IdxNum = getNumElements(),
1473 typename = EnableIfOneIndex<IdxNum>,
1474 typename = EnableIfScalarType<T>>
1475 operator T() const {
1479 template <typename T, typename = EnableIfScalarType<T>>
1480 friend NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>
1481 operator*(const T &Lhs, const SwizzleOp &Rhs) {
1482 return NewRHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
1483 Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1486 template <typename T, typename = EnableIfScalarType<T>>
1487 friend NewRHOp<GetScalarOp<T>, std::plus, Indexes...>
1488 operator+(const T &Lhs, const SwizzleOp &Rhs) {
1489 return NewRHOp<GetScalarOp<T>, std::plus, Indexes...>(
1490 Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1493 template <typename T, typename = EnableIfScalarType<T>>
1494 friend NewRHOp<GetScalarOp<T>, std::divides, Indexes...>
1495 operator/(const T &Lhs, const SwizzleOp &Rhs) {
1496 return NewRHOp<GetScalarOp<T>, std::divides, Indexes...>(
1497 Rhs.m_Vector, GetScalarOp<T>(Lhs), Rhs);
1501 #ifdef __SYCL_OPASSIGN
1502 #error "Undefine __SYCL_OPASSIGN macro."
1504 #define __SYCL_OPASSIGN(OPASSIGN, OP) \
1505 SwizzleOp &operator OPASSIGN(const DataT &Rhs) { \
1506 operatorHelper<OP>(vec_t(Rhs)); \
1509 template <typename RhsOperation> \
1510 SwizzleOp &operator OPASSIGN(const RhsOperation &Rhs) { \
1511 operatorHelper<OP>(Rhs); \
1515 __SYCL_OPASSIGN(+=, std::plus)
1516 __SYCL_OPASSIGN(-=, std::minus)
1517 __SYCL_OPASSIGN(*=, std::multiplies)
1518 __SYCL_OPASSIGN(/=, std::divides)
1519 __SYCL_OPASSIGN(%=, std::modulus)
1520 __SYCL_OPASSIGN(&=, std::bit_and)
1521 __SYCL_OPASSIGN(|=, std::bit_or)
1522 __SYCL_OPASSIGN(^=, std::bit_xor)
1523 __SYCL_OPASSIGN(>>=, RShift)
1524 __SYCL_OPASSIGN(<<=, LShift)
1525 #undef __SYCL_OPASSIGN
1528 #error "Undefine __SYCL_UOP macro"
1530 #define __SYCL_UOP(UOP, OPASSIGN) \
1531 SwizzleOp &operator UOP() { \
1532 *this OPASSIGN static_cast<DataT>(1); \
1535 vec_t operator UOP(int) { \
1536 vec_t Ret = *this; \
1537 *this OPASSIGN static_cast<DataT>(1); \
1545 template <typename T = DataT>
1546 typename detail::enable_if_t<std::is_integral<vec_data_t<T>>::value, vec_t>
1552 vec_rel_t operator!() {
1567 template <int IdxNum = getNumElements(),
1568 typename = EnableIfMultipleIndexes<IdxNum>>
1569 SwizzleOp &operator=(const vec<DataT, IdxNum> &Rhs) {
1570 std::array<int, IdxNum> Idxs{Indexes...};
1571 for (size_t I = 0; I < Idxs.size(); ++I) {
1572 m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1577 template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
1578 SwizzleOp &operator=(const DataT &Rhs) {
1579 std::array<int, IdxNum> Idxs{Indexes...};
1580 m_Vector->setValue(Idxs[0], Rhs);
1584 template <int IdxNum = getNumElements(), typename = EnableIfOneIndex<IdxNum>>
1585 SwizzleOp &operator=(DataT &&Rhs) {
1586 std::array<int, IdxNum> Idxs{Indexes...};
1587 m_Vector->setValue(Idxs[0], Rhs);
1591 template <typename T, typename = EnableIfScalarType<T>>
1592 NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>
1593 operator*(const T &Rhs) const {
1594 return NewLHOp<GetScalarOp<T>, std::multiplies, Indexes...>(
1595 m_Vector, *this, GetScalarOp<T>(Rhs));
1598 template <typename RhsOperation,
1599 typename = EnableIfNoScalarType<RhsOperation>>
1600 NewLHOp<RhsOperation, std::multiplies, Indexes...>
1601 operator*(const RhsOperation &Rhs) const {
1602 return NewLHOp<RhsOperation, std::multiplies, Indexes...>(m_Vector, *this,
1606 template <typename T, typename = EnableIfScalarType<T>>
1607 NewLHOp<GetScalarOp<T>, std::plus, Indexes...> operator+(const T &Rhs) const {
1608 return NewLHOp<GetScalarOp<T>, std::plus, Indexes...>(m_Vector, *this,
1609 GetScalarOp<T>(Rhs));
1612 template <typename RhsOperation,
1613 typename = EnableIfNoScalarType<RhsOperation>>
1614 NewLHOp<RhsOperation, std::plus, Indexes...>
1615 operator+(const RhsOperation &Rhs) const {
1616 return NewLHOp<RhsOperation, std::plus, Indexes...>(m_Vector, *this, Rhs);
1619 template <typename T, typename = EnableIfScalarType<T>>
1620 NewLHOp<GetScalarOp<T>, std::minus, Indexes...>
1621 operator-(const T &Rhs) const {
1622 return NewLHOp<GetScalarOp<T>, std::minus, Indexes...>(m_Vector, *this,
1623 GetScalarOp<T>(Rhs));
1626 template <typename RhsOperation,
1627 typename = EnableIfNoScalarType<RhsOperation>>
1628 NewLHOp<RhsOperation, std::minus, Indexes...>
1629 operator-(const RhsOperation &Rhs) const {
1630 return NewLHOp<RhsOperation, std::minus, Indexes...>(m_Vector, *this, Rhs);
1633 template <typename T, typename = EnableIfScalarType<T>>
1634 NewLHOp<GetScalarOp<T>, std::divides, Indexes...>
1635 operator/(const T &Rhs) const {
1636 return NewLHOp<GetScalarOp<T>, std::divides, Indexes...>(
1637 m_Vector, *this, GetScalarOp<T>(Rhs));
1640 template <typename RhsOperation,
1641 typename = EnableIfNoScalarType<RhsOperation>>
1642 NewLHOp<RhsOperation, std::divides, Indexes...>
1643 operator/(const RhsOperation &Rhs) const {
1644 return NewLHOp<RhsOperation, std::divides, Indexes...>(m_Vector, *this,
1648 template <typename T, typename = EnableIfScalarType<T>>
1649 NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>
1650 operator&(const T &Rhs) const {
1651 return NewLHOp<GetScalarOp<T>, std::bit_and, Indexes...>(
1652 m_Vector, *this, GetScalarOp<T>(Rhs));
1655 template <typename RhsOperation,
1656 typename = EnableIfNoScalarType<RhsOperation>>
1657 NewLHOp<RhsOperation, std::bit_and, Indexes...>
1658 operator&(const RhsOperation &Rhs) const {
1659 return NewLHOp<RhsOperation, std::bit_and, Indexes...>(m_Vector, *this,
1663 template <typename T, typename = EnableIfScalarType<T>>
1664 NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>
1665 operator|(const T &Rhs) const {
1666 return NewLHOp<GetScalarOp<T>, std::bit_or, Indexes...>(
1667 m_Vector, *this, GetScalarOp<T>(Rhs));
1670 template <typename RhsOperation,
1671 typename = EnableIfNoScalarType<RhsOperation>>
1672 NewLHOp<RhsOperation, std::bit_or, Indexes...>
1673 operator|(const RhsOperation &Rhs) const {
1674 return NewLHOp<RhsOperation, std::bit_or, Indexes...>(m_Vector, *this, Rhs);
1677 template <typename T, typename = EnableIfScalarType<T>>
1678 NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>
1679 operator^(const T &Rhs) const {
1680 return NewLHOp<GetScalarOp<T>, std::bit_xor, Indexes...>(
1681 m_Vector, *this, GetScalarOp<T>(Rhs));
1684 template <typename RhsOperation,
1685 typename = EnableIfNoScalarType<RhsOperation>>
1686 NewLHOp<RhsOperation, std::bit_xor, Indexes...>
1687 operator^(const RhsOperation &Rhs) const {
1688 return NewLHOp<RhsOperation, std::bit_xor, Indexes...>(m_Vector, *this,
1692 template <typename T, typename = EnableIfScalarType<T>>
1693 NewLHOp<GetScalarOp<T>, RShift, Indexes...> operator>>(const T &Rhs) const {
1694 return NewLHOp<GetScalarOp<T>, RShift, Indexes...>(m_Vector, *this,
1695 GetScalarOp<T>(Rhs));
1698 template <typename RhsOperation,
1699 typename = EnableIfNoScalarType<RhsOperation>>
1700 NewLHOp<RhsOperation, RShift, Indexes...>
1701 operator>>(const RhsOperation &Rhs) const {
1702 return NewLHOp<RhsOperation, RShift, Indexes...>(m_Vector, *this, Rhs);
1705 template <typename T, typename = EnableIfScalarType<T>>
1706 NewLHOp<GetScalarOp<T>, LShift, Indexes...> operator<<(const T &Rhs) const {
1707 return NewLHOp<GetScalarOp<T>, LShift, Indexes...>(m_Vector, *this,
1708 GetScalarOp<T>(Rhs));
1711 template <typename RhsOperation,
1712 typename = EnableIfNoScalarType<RhsOperation>>
1713 NewLHOp<RhsOperation, LShift, Indexes...>
1714 operator<<(const RhsOperation &Rhs) const {
1715 return NewLHOp<RhsOperation, LShift, Indexes...>(m_Vector, *this, Rhs);
1718 template <typename T1, typename T2, typename T3, template <typename> class T4,
1721 typename detail::enable_if_t<sizeof...(T5) == getNumElements()>>
1722 SwizzleOp &operator=(const SwizzleOp<T1, T2, T3, T4, T5...> &Rhs) {
1723 std::array<int, getNumElements()> Idxs{Indexes...};
1724 for (size_t I = 0; I < Idxs.size(); ++I) {
1725 m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1730 template <typename T1, typename T2, typename T3, template <typename> class T4,
1733 typename detail::enable_if_t<sizeof...(T5) == getNumElements()>>
1734 SwizzleOp &operator=(SwizzleOp<T1, T2, T3, T4, T5...> &&Rhs) {
1735 std::array<int, getNumElements()> Idxs{Indexes...};
1736 for (size_t I = 0; I < Idxs.size(); ++I) {
1737 m_Vector->setValue(Idxs[I], Rhs.getValue(I));
1742 template <typename T, typename = EnableIfScalarType<T>>
1743 NewRelOp<GetScalarOp<T>, EqualTo, Indexes...> operator==(const T &Rhs) const {
1744 return NewRelOp<GetScalarOp<T>, EqualTo, Indexes...>(NULL, *this,
1745 GetScalarOp<T>(Rhs));
1748 template <typename RhsOperation,
1749 typename = EnableIfNoScalarType<RhsOperation>>
1750 NewRelOp<RhsOperation, EqualTo, Indexes...>
1751 operator==(const RhsOperation &Rhs) const {
1752 return NewRelOp<RhsOperation, EqualTo, Indexes...>(NULL, *this, Rhs);
1755 template <typename T, typename = EnableIfScalarType<T>>
1756 NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>
1757 operator!=(const T &Rhs) const {
1758 return NewRelOp<GetScalarOp<T>, NotEqualTo, Indexes...>(
1759 NULL, *this, GetScalarOp<T>(Rhs));
1762 template <typename RhsOperation,
1763 typename = EnableIfNoScalarType<RhsOperation>>
1764 NewRelOp<RhsOperation, NotEqualTo, Indexes...>
1765 operator!=(const RhsOperation &Rhs) const {
1766 return NewRelOp<RhsOperation, NotEqualTo, Indexes...>(NULL, *this, Rhs);
1769 template <typename T, typename = EnableIfScalarType<T>>
1770 NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>
1771 operator>=(const T &Rhs) const {
1772 return NewRelOp<GetScalarOp<T>, GreaterEqualTo, Indexes...>(
1773 NULL, *this, GetScalarOp<T>(Rhs));
1776 template <typename RhsOperation,
1777 typename = EnableIfNoScalarType<RhsOperation>>
1778 NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>
1779 operator>=(const RhsOperation &Rhs) const {
1780 return NewRelOp<RhsOperation, GreaterEqualTo, Indexes...>(NULL, *this, Rhs);
1783 template <typename T, typename = EnableIfScalarType<T>>
1784 NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>
1785 operator<=(const T &Rhs) const {
1786 return NewRelOp<GetScalarOp<T>, LessEqualTo, Indexes...>(
1787 NULL, *this, GetScalarOp<T>(Rhs));
1790 template <typename RhsOperation,
1791 typename = EnableIfNoScalarType<RhsOperation>>
1792 NewRelOp<RhsOperation, LessEqualTo, Indexes...>
1793 operator<=(const RhsOperation &Rhs) const {
1794 return NewRelOp<RhsOperation, LessEqualTo, Indexes...>(NULL, *this, Rhs);
1797 template <typename T, typename = EnableIfScalarType<T>>
1798 NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>
1799 operator>(const T &Rhs) const {
1800 return NewRelOp<GetScalarOp<T>, GreaterThan, Indexes...>(
1801 NULL, *this, GetScalarOp<T>(Rhs));
1804 template <typename RhsOperation,
1805 typename = EnableIfNoScalarType<RhsOperation>>
1806 NewRelOp<RhsOperation, GreaterThan, Indexes...>
1807 operator>(const RhsOperation &Rhs) const {
1808 return NewRelOp<RhsOperation, GreaterThan, Indexes...>(NULL, *this, Rhs);
1811 template <typename T, typename = EnableIfScalarType<T>>
1812 NewRelOp<GetScalarOp<T>, LessThan, Indexes...> operator<(const T &Rhs) const {
1813 return NewRelOp<GetScalarOp<T>, LessThan, Indexes...>(NULL, *this,
1814 GetScalarOp<T>(Rhs));
1817 template <typename RhsOperation,
1818 typename = EnableIfNoScalarType<RhsOperation>>
1819 NewRelOp<RhsOperation, LessThan, Indexes...>
1820 operator<(const RhsOperation &Rhs) const {
1821 return NewRelOp<RhsOperation, LessThan, Indexes...>(NULL, *this, Rhs);
1824 template <typename T, typename = EnableIfScalarType<T>>
1825 NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>
1826 operator&&(const T &Rhs) const {
1827 return NewRelOp<GetScalarOp<T>, LogicalAnd, Indexes...>(
1828 NULL, *this, GetScalarOp<T>(Rhs));
1831 template <typename RhsOperation,
1832 typename = EnableIfNoScalarType<RhsOperation>>
1833 NewRelOp<RhsOperation, LogicalAnd, Indexes...>
1834 operator&&(const RhsOperation &Rhs) const {
1835 return NewRelOp<RhsOperation, LogicalAnd, Indexes...>(NULL, *this, Rhs);
1838 template <typename T, typename = EnableIfScalarType<T>>
1839 NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>
1840 operator||(const T &Rhs) const {
1841 return NewRelOp<GetScalarOp<T>, LogicalOr, Indexes...>(NULL, *this,
1842 GetScalarOp<T>(Rhs));
1845 template <typename RhsOperation,
1846 typename = EnableIfNoScalarType<RhsOperation>>
1847 NewRelOp<RhsOperation, LogicalOr, Indexes...>
1848 operator||(const RhsOperation &Rhs) const {
1849 return NewRelOp<RhsOperation, LogicalOr, Indexes...>(NULL, *this, Rhs);
1859 template <int Index>
1861 static constexpr int IDXs[sizeof...(Indexes)] = {Indexes...};
1862 static constexpr int value = IDXs[Index >= getNumElements() ? 0 : Index];
1866 #ifdef __SYCL_ACCESS_RETURN
1867 #error "Undefine __SYCL_ACCESS_RETURN macro"
1869 #define __SYCL_ACCESS_RETURN m_Vector
1870 #include "swizzles.def"
1871 #undef __SYCL_ACCESS_RETURN
1876 template <access::address_space Space>
1877 void load(size_t offset, multi_ptr<DataT, Space> ptr) {
1879 Tmp.template load(offset, ptr);
1883 template <typename convertT, rounding_mode roundingMode>
1884 vec<convertT, sizeof...(Indexes)> convert() const {
1887 return Tmp.template convert<convertT, roundingMode>();
1890 template <typename asT> asT as() const {
1893 static_assert((sizeof(Tmp) == sizeof(asT)),
1894 "The new SYCL vec type must have the same storage size in "
1895 "bytes as this SYCL swizzled vec");
1897 detail::is_contained<asT, detail::gtl::vector_basic_list>::value,
1898 "asT must be SYCL vec of a different element type and "
1899 "number of elements specified by asT");
1900 return Tmp.template as<asT>();
1904 SwizzleOp(const SwizzleOp &Rhs)
1905 : m_Vector(Rhs.m_Vector), m_LeftOperation(Rhs.m_LeftOperation),
1906 m_RightOperation(Rhs.m_RightOperation) {}
1908 SwizzleOp(VecT *Vector, OperationLeftT LeftOperation,
1909 OperationRightT RightOperation)
1910 : m_Vector(Vector), m_LeftOperation(LeftOperation),
1911 m_RightOperation(RightOperation) {}
1913 SwizzleOp(VecT *Vector) : m_Vector(Vector) {}
1915 SwizzleOp(SwizzleOp &&Rhs)
1916 : m_Vector(Rhs.m_Vector), m_LeftOperation(std::move(Rhs.m_LeftOperation)),
1917 m_RightOperation(std::move(Rhs.m_RightOperation)) {}
1923 template <int IdxNum = getNumElements()>
1924 CommonDataT getValue(EnableIfOneIndex<IdxNum, size_t> Index) const {
1925 if (std::is_same<OperationCurrentT<DataT>, GetOp<DataT>>::value) {
1926 std::array<int, getNumElements()> Idxs{Indexes...};
1927 return m_Vector->getValue(Idxs[Index]);
1929 auto Op = OperationCurrentT<vec_data_t<CommonDataT>>();
1930 return vec_data<CommonDataT>::get(
1931 Op(vec_data<CommonDataT>::get(m_LeftOperation.getValue(Index)),
1932 vec_data<CommonDataT>::get(m_RightOperation.getValue(Index))));
1935 template <int IdxNum = getNumElements()>
1936 DataT getValue(EnableIfMultipleIndexes<IdxNum, size_t> Index) const {
1937 if (std::is_same<OperationCurrentT<DataT>, GetOp<DataT>>::value) {
1938 std::array<int, getNumElements()> Idxs{Indexes...};
1939 return m_Vector->getValue(Idxs[Index]);
1941 auto Op = OperationCurrentT<vec_data_t<DataT>>();
1942 return vec_data<DataT>::get(
1943 Op(vec_data<DataT>::get(m_LeftOperation.getValue(Index)),
1944 vec_data<DataT>::get(m_RightOperation.getValue(Index))));
1947 template <template <typename> class Operation, typename RhsOperation>
1948 void operatorHelper(const RhsOperation &Rhs) {
1949 Operation<vec_data_t<DataT>> Op;
1950 std::array<int, getNumElements()> Idxs{Indexes...};
1951 for (size_t I = 0; I < Idxs.size(); ++I) {
1952 DataT Res = vec_data<DataT>::get(
1953 Op(vec_data<DataT>::get(m_Vector->getValue(Idxs[I])),
1954 vec_data<DataT>::get(Rhs.getValue(I))));
1955 m_Vector->setValue(Idxs[I], Res);
1962 OperationLeftT m_LeftOperation;
1963 OperationRightT m_RightOperation;
1966 template <typename T1, int T2> friend class cl::sycl::vec;
1968 template <typename T1, typename T2, typename T3, template <typename> class T4,
1970 friend class SwizzleOp;
1978 #error "Undefine __SYCL_BINOP macro"
1980 #define __SYCL_BINOP(BINOP) \
1981 template <typename T, int Num> \
1982 typename detail::enable_if_t< \
1983 std::is_fundamental<vec_data_t<T>>::value || \
1984 std::is_same<typename detail::remove_const_t<T>, half>::value, \
1986 operator BINOP(const T &Lhs, const vec<T, Num> &Rhs) { \
1987 return vec<T, Num>(Lhs) BINOP Rhs; \
1989 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
1990 template <typename> class OperationCurrentT, int... Indexes, \
1991 typename T, typename T1 = typename VecT::element_type, \
1992 int Num = sizeof...(Indexes)> \
1993 typename detail::enable_if_t< \
1994 std::is_convertible<T, T1>::value && \
1995 (std::is_fundamental<vec_data_t<T>>::value || \
1996 std::is_same<typename detail::remove_const_t<T>, half>::value), \
2000 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2001 OperationCurrentT, Indexes...> &Rhs) { \
2002 vec<T1, Num> Tmp = Rhs; \
2003 return Lhs BINOP Tmp; \
2005 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2006 template <typename> class OperationCurrentT, int... Indexes, \
2007 typename T = typename VecT::element_type, \
2008 int Num = sizeof...(Indexes)> \
2009 vec<T, Num> operator BINOP( \
2010 const vec<T, Num> &Lhs, \
2011 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2012 OperationCurrentT, Indexes...> &Rhs) { \
2013 vec<T, Num> Tmp = Rhs; \
2014 return Lhs BINOP Tmp; \
2031 #ifdef __SYCL_RELLOGOP
2032 #error "Undefine __SYCL_RELLOGOP macro"
2034 #define __SYCL_RELLOGOP(RELLOGOP) \
2035 template <typename T, typename DataT, int Num> \
2036 typename detail::enable_if_t< \
2037 std::is_convertible<T, DataT>::value && \
2038 (std::is_fundamental<vec_data_t<T>>::value || \
2039 std::is_same<typename detail::remove_const_t<T>, half>::value), \
2040 vec<detail::rel_t<DataT>, Num>> \
2041 operator RELLOGOP(const T &Lhs, const vec<DataT, Num> &Rhs) { \
2042 return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs; \
2044 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2045 template <typename> class OperationCurrentT, int... Indexes, \
2046 typename T, typename T1 = typename VecT::element_type, \
2047 int Num = sizeof...(Indexes)> \
2048 typename detail::enable_if_t< \
2049 std::is_convertible<T, T1>::value && \
2050 (std::is_fundamental<vec_data_t<T>>::value || \
2051 std::is_same<typename detail::remove_const_t<T>, half>::value), \
2052 vec<detail::rel_t<T1>, Num>> \
2053 operator RELLOGOP( \
2055 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2056 OperationCurrentT, Indexes...> &Rhs) { \
2057 vec<T1, Num> Tmp = Rhs; \
2058 return Lhs RELLOGOP Tmp; \
2060 template <typename VecT, typename OperationLeftT, typename OperationRightT, \
2061 template <typename> class OperationCurrentT, int... Indexes, \
2062 typename T = typename VecT::element_type, \
2063 int Num = sizeof...(Indexes)> \
2064 vec<detail::rel_t<T>, Num> operator RELLOGOP( \
2065 const vec<T, Num> &Lhs, \
2066 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT, \
2067 OperationCurrentT, Indexes...> &Rhs) { \
2068 vec<T, Num> Tmp = Rhs; \
2069 return Lhs RELLOGOP Tmp; \
2081 #undef __SYCL_RELLOGOP
2087 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
2088 #define __SYCL_DECLARE_TYPE_VIA_CL_T(type) \
2089 using __##type##_t = cl::sycl::cl_##type; \
2090 using __##type##2_vec_t = \
2091 cl::sycl::cl_##type __attribute__((ext_vector_type(2))); \
2092 using __##type##3_vec_t = \
2093 cl::sycl::cl_##type __attribute__((ext_vector_type(3))); \
2094 using __##type##4_vec_t = \
2095 cl::sycl::cl_##type __attribute__((ext_vector_type(4))); \
2096 using __##type##8_vec_t = \
2097 cl::sycl::cl_##type __attribute__((ext_vector_type(8))); \
2098 using __##type##16_vec_t = \
2099 cl::sycl::cl_##type __attribute__((ext_vector_type(16)));
2101 #define __SYCL_DECLARE_TYPE_T(type) \
2102 using __##type##_t = cl::sycl::type; \
2103 using __##type##2_vec_t = \
2104 cl::sycl::type __attribute__((ext_vector_type(2))); \
2105 using __##type##3_vec_t = \
2106 cl::sycl::type __attribute__((ext_vector_type(3))); \
2107 using __##type##4_vec_t = \
2108 cl::sycl::type __attribute__((ext_vector_type(4))); \
2109 using __##type##8_vec_t = \
2110 cl::sycl::type __attribute__((ext_vector_type(8))); \
2111 using __##type##16_vec_t = \
2112 cl::sycl::type __attribute__((ext_vector_type(16)));
2114 __SYCL_DECLARE_TYPE_VIA_CL_T(char)
2115 __SYCL_DECLARE_TYPE_T(schar)
2116 __SYCL_DECLARE_TYPE_VIA_CL_T(uchar)
2117 __SYCL_DECLARE_TYPE_VIA_CL_T(short)
2118 __SYCL_DECLARE_TYPE_VIA_CL_T(ushort)
2119 __SYCL_DECLARE_TYPE_VIA_CL_T(int)
2120 __SYCL_DECLARE_TYPE_VIA_CL_T(uint)
2121 __SYCL_DECLARE_TYPE_VIA_CL_T(long)
2122 __SYCL_DECLARE_TYPE_VIA_CL_T(ulong)
2123 __SYCL_DECLARE_TYPE_T(longlong)
2124 __SYCL_DECLARE_TYPE_T(ulonglong)
2127 __SYCL_DECLARE_TYPE_VIA_CL_T(float)
2128 __SYCL_DECLARE_TYPE_VIA_CL_T(double)
2130 #define __SYCL_GET_CL_TYPE(target, num) __##target##num##_vec_t
2131 #define __SYCL_GET_SCALAR_CL_TYPE(target) target
2133 #undef __SYCL_DECLARE_TYPE_VIA_CL_T
2134 #undef __SYCL_DECLARE_TYPE_T
2136 #define __SYCL_GET_CL_TYPE(target, num) ::cl_##target##num
2137 #define __SYCL_GET_SCALAR_CL_TYPE(target) ::cl_##target
2140 using __half_t = cl::sycl::detail::half_impl::StorageT;
2141 using __half2_vec_t = cl::sycl::detail::half_impl::Vec2StorageT;
2142 using __half3_vec_t = cl::sycl::detail::half_impl::Vec3StorageT;
2143 using __half4_vec_t = cl::sycl::detail::half_impl::Vec4StorageT;
2144 using __half8_vec_t = cl::sycl::detail::half_impl::Vec8StorageT;
2145 using __half16_vec_t = cl::sycl::detail::half_impl::Vec16StorageT;
2146 #define __SYCL_GET_CL_HALF_TYPE(target, num) __##target##num##_vec_t
2148 __SYCL_INLINE_NAMESPACE(cl) {
2153 template <typename T, typename T8, typename T16, typename T32, typename T64>
2154 using select_apply_cl_t =
2155 conditional_t<sizeof(T) == 1, T8,
2156 conditional_t<sizeof(T) == 2, T16,
2157 conditional_t<sizeof(T) == 4, T32, T64>>>;
2160 #define __SYCL_DECLARE_CONVERTER(base, num) \
2161 template <> class BaseCLTypeConverter<base, num> { \
2163 using DataType = __SYCL_GET_CL_TYPE(base, num); \
2166 #define __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, num) \
2167 template <> class BaseCLTypeConverter<base, num> { \
2169 using DataType = detail::select_apply_cl_t< \
2170 base, __SYCL_GET_CL_TYPE(char, num), __SYCL_GET_CL_TYPE(short, num), \
2171 __SYCL_GET_CL_TYPE(int, num), __SYCL_GET_CL_TYPE(long, num)>; \
2174 #define __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, num) \
2175 template <> class BaseCLTypeConverter<base, num> { \
2177 using DataType = detail::select_apply_cl_t< \
2178 base, __SYCL_GET_CL_TYPE(uchar, num), __SYCL_GET_CL_TYPE(ushort, num), \
2179 __SYCL_GET_CL_TYPE(uint, num), __SYCL_GET_CL_TYPE(ulong, num)>; \
2182 #define __SYCL_DECLARE_FLOAT_CONVERTER(base, num) \
2183 template <> class BaseCLTypeConverter<base, num> { \
2185 using DataType = detail::select_apply_cl_t< \
2186 base, std::false_type, __SYCL_GET_CL_HALF_TYPE(half, num), \
2187 __SYCL_GET_CL_TYPE(float, num), __SYCL_GET_CL_TYPE(double, num)>; \
2190 #define __SYCL_DECLARE_LONGLONG_CONVERTER(base, num) \
2191 template <> class BaseCLTypeConverter<base##long, num> { \
2193 using DataType = __SYCL_GET_CL_TYPE(base, num); \
2196 #define __SYCL_DECLARE_SCHAR_CONVERTER(num) \
2197 template <> class BaseCLTypeConverter<schar, num> { \
2199 using DataType = detail::select_apply_cl_t< \
2200 schar, __SYCL_GET_CL_TYPE(char, num), __SYCL_GET_CL_TYPE(short, num), \
2201 __SYCL_GET_CL_TYPE(int, num), __SYCL_GET_CL_TYPE(long, num)>; \
2204 #define __SYCL_DECLARE_BOOL_CONVERTER(num) \
2205 template <> class BaseCLTypeConverter<bool, num> { \
2207 using DataType = detail::select_apply_cl_t< \
2208 bool, __SYCL_GET_CL_TYPE(char, num), __SYCL_GET_CL_TYPE(short, num), \
2209 __SYCL_GET_CL_TYPE(int, num), __SYCL_GET_CL_TYPE(long, num)>; \
2212 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2213 #define __SYCL_DECLARE_BYTE_CONVERTER(num) \
2214 template <> class BaseCLTypeConverter<std::byte, num> { \
2216 using DataType = __SYCL_GET_CL_TYPE(uchar, num); \
2219 #define __SYCL_DECLARE_HALF_CONVERTER(base, num) \
2220 template <> class BaseCLTypeConverter<base, num> { \
2222 using DataType = __SYCL_GET_CL_HALF_TYPE(base, num); \
2225 #define __SYCL_DECLARE_SCALAR_SCHAR_CONVERTER \
2226 template <> class BaseCLTypeConverter<schar, 1> { \
2228 using DataType = schar; \
2231 #define __SYCL_DECLARE_SCALAR_BOOL_CONVERTER \
2232 template <> class BaseCLTypeConverter<bool, 1> { \
2234 using DataType = bool; \
2237 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2238 #define __SYCL_DECLARE_SCALAR_BYTE_CONVERTER \
2239 template <> class BaseCLTypeConverter<std::byte, 1> { \
2241 using DataType = __SYCL_GET_SCALAR_CL_TYPE(uchar); \
2244 #define __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2245 template <> class BaseCLTypeConverter<base, 1> { \
2247 using DataType = __SYCL_GET_SCALAR_CL_TYPE(base); \
2250 #define __SYCL_DECLARE_VECTOR_CONVERTERS(base) \
2251 namespace detail { \
2252 __SYCL_DECLARE_CONVERTER(base, 2) \
2253 __SYCL_DECLARE_CONVERTER(base, 3) \
2254 __SYCL_DECLARE_CONVERTER(base, 4) \
2255 __SYCL_DECLARE_CONVERTER(base, 8) \
2256 __SYCL_DECLARE_CONVERTER(base, 16) \
2257 __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2260 #define __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(base) \
2261 namespace detail { \
2262 __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 2) \
2263 __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 3) \
2264 __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 4) \
2265 __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 8) \
2266 __SYCL_DECLARE_SIGNED_INTEGRAL_CONVERTER(base, 16) \
2267 __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2270 #define __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(base) \
2271 namespace detail { \
2272 __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 2) \
2273 __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 3) \
2274 __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 4) \
2275 __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 8) \
2276 __SYCL_DECLARE_UNSIGNED_INTEGRAL_CONVERTER(base, 16) \
2277 __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2280 #define __SYCL_DECLARE_FLOAT_VECTOR_CONVERTERS(base) \
2281 namespace detail { \
2282 __SYCL_DECLARE_FLOAT_CONVERTER(base, 2) \
2283 __SYCL_DECLARE_FLOAT_CONVERTER(base, 3) \
2284 __SYCL_DECLARE_FLOAT_CONVERTER(base, 4) \
2285 __SYCL_DECLARE_FLOAT_CONVERTER(base, 8) \
2286 __SYCL_DECLARE_FLOAT_CONVERTER(base, 16) \
2287 __SYCL_DECLARE_SCALAR_CONVERTER(base) \
2290 #define __SYCL_DECLARE_HALF_VECTOR_CONVERTERS(base) \
2291 namespace detail { \
2292 __SYCL_DECLARE_HALF_CONVERTER(base, 2) \
2293 __SYCL_DECLARE_HALF_CONVERTER(base, 3) \
2294 __SYCL_DECLARE_HALF_CONVERTER(base, 4) \
2295 __SYCL_DECLARE_HALF_CONVERTER(base, 8) \
2296 __SYCL_DECLARE_HALF_CONVERTER(base, 16) \
2297 template <> class BaseCLTypeConverter<base, 1> { \
2299 using DataType = __half_t; \
2303 #define __SYCL_DECLARE_VECTOR_LONGLONG_CONVERTERS(base) \
2304 namespace detail { \
2305 __SYCL_DECLARE_LONGLONG_CONVERTER(base, 2) \
2306 __SYCL_DECLARE_LONGLONG_CONVERTER(base, 3) \
2307 __SYCL_DECLARE_LONGLONG_CONVERTER(base, 4) \
2308 __SYCL_DECLARE_LONGLONG_CONVERTER(base, 8) \
2309 __SYCL_DECLARE_LONGLONG_CONVERTER(base, 16) \
2310 template <> class BaseCLTypeConverter<base##long, 1> { \
2312 using DataType = base##long; \
2316 #define __SYCL_DECLARE_SCHAR_VECTOR_CONVERTERS \
2317 namespace detail { \
2318 __SYCL_DECLARE_SCHAR_CONVERTER(2) \
2319 __SYCL_DECLARE_SCHAR_CONVERTER(3) \
2320 __SYCL_DECLARE_SCHAR_CONVERTER(4) \
2321 __SYCL_DECLARE_SCHAR_CONVERTER(8) \
2322 __SYCL_DECLARE_SCHAR_CONVERTER(16) \
2323 __SYCL_DECLARE_SCALAR_SCHAR_CONVERTER \
2326 #define __SYCL_DECLARE_BOOL_VECTOR_CONVERTERS \
2327 namespace detail { \
2328 __SYCL_DECLARE_BOOL_CONVERTER(2) \
2329 __SYCL_DECLARE_BOOL_CONVERTER(3) \
2330 __SYCL_DECLARE_BOOL_CONVERTER(4) \
2331 __SYCL_DECLARE_BOOL_CONVERTER(8) \
2332 __SYCL_DECLARE_BOOL_CONVERTER(16) \
2333 __SYCL_DECLARE_SCALAR_BOOL_CONVERTER \
2336 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2337 #define __SYCL_DECLARE_BYTE_VECTOR_CONVERTER \
2338 namespace detail { \
2339 __SYCL_DECLARE_BYTE_CONVERTER(2) \
2340 __SYCL_DECLARE_BYTE_CONVERTER(3) \
2341 __SYCL_DECLARE_BYTE_CONVERTER(4) \
2342 __SYCL_DECLARE_BYTE_CONVERTER(8) \
2343 __SYCL_DECLARE_BYTE_CONVERTER(16) \
2344 __SYCL_DECLARE_SCALAR_BYTE_CONVERTER \
2347 __SYCL_DECLARE_VECTOR_CONVERTERS(char)
2348 __SYCL_DECLARE_SCHAR_VECTOR_CONVERTERS
2349 __SYCL_DECLARE_BOOL_VECTOR_CONVERTERS
2350 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2351 __SYCL_DECLARE_BYTE_VECTOR_CONVERTER
2353 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(uchar)
2354 __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(short)
2355 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(ushort)
2356 __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(int)
2357 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(uint)
2358 __SYCL_DECLARE_SIGNED_INTEGRAL_VECTOR_CONVERTERS(long)
2359 __SYCL_DECLARE_UNSIGNED_INTEGRAL_VECTOR_CONVERTERS(ulong)
2360 __SYCL_DECLARE_VECTOR_LONGLONG_CONVERTERS(long)
2361 __SYCL_DECLARE_VECTOR_LONGLONG_CONVERTERS(ulong)
2362 __SYCL_DECLARE_HALF_VECTOR_CONVERTERS(half)
2363 __SYCL_DECLARE_FLOAT_VECTOR_CONVERTERS(float)
2364 __SYCL_DECLARE_FLOAT_VECTOR_CONVERTERS(double)
2366 #undef __SYCL_GET_CL_TYPE
2367 #undef __SYCL_GET_SCALAR_CL_TYPE
2368 #undef __SYCL_DECLARE_CONVERTER
2369 #undef __SYCL_DECLARE_VECTOR_CONVERTERS
2370 #undef __SYCL_DECLARE_SYCL_VEC
2371 #undef __SYCL_DECLARE_SYCL_VEC_WO_CONVERTERS
2372 #undef __SYCL_DECLARE_SCHAR_VECTOR_CONVERTERS
2373 #undef __SYCL_DECLARE_SCHAR_CONVERTER
2374 #undef __SYCL_DECLARE_SCALAR_SCHAR_CONVERTER
2375 #undef __SYCL_DECLARE_BOOL_VECTOR_CONVERTERS
2376 #undef __SYCL_DECLARE_BOOL_CONVERTER
2377 #if __cplusplus >= 201703L && (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2378 #undef __SYCL_DECLARE_BYTE_VECTOR_CONVERTER
2379 #undef __SYCL_DECLARE_BYTE_CONVERTER
2380 #undef __SYCL_DECLARE_SCALAR_BYTE_CONVERTER
2382 #undef __SYCL_DECLARE_SCALAR_BOOL_CONVERTER
2383 #undef __SYCL_USE_EXT_VECTOR_TYPE__
2388 #define SYCL_DEVICE_COPYABLE 1
2397 template <typename T, typename = void>
2398 struct is_device_copyable : std::false_type {};
2400 template <typename T>
2401 struct is_device_copyable<
2402 T, std::enable_if_t<std::is_trivially_copyable<T>::value>>
2403 : std::true_type {};
2405 #if __cplusplus >= 201703L
2406 template <typename T>
2407 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
2411 template <> struct is_device_copyable<std::tuple<>> : std::true_type {};
2415 template <typename T, typename... Ts>
2416 struct is_device_copyable<std::tuple<T, Ts...>>
2417 : detail::bool_constant<is_device_copyable<T>::value &&
2418 is_device_copyable<std::tuple<Ts...>>::value> {};
2423 template <typename T, std::size_t N>
2424 struct is_device_copyable<
2425 sycl::marray<T, N>, std::enable_if_t<is_device_copyable<T>::value &&
2426 !std::is_trivially_copyable<T>::value>>
2427 : std::true_type {};
2430 template <typename T, typename = void>
2431 struct IsDeprecatedDeviceCopyable : std::false_type {};
2435 template <typename T>
2436 struct __SYCL2020_DEPRECATED("This type isn't device copyable in SYCL 2020")
2437 IsDeprecatedDeviceCopyable<
2438 T, std::enable_if_t<std::is_trivially_copy_constructible<T>::value &&
2439 std::is_trivially_destructible<T>::value &&
2440 !is_device_copyable<T>::value>> : std::true_type {};
2442 #ifdef __SYCL_DEVICE_ONLY__
2445 template <typename T, unsigned NumFieldsToCheck>
2446 struct CheckFieldsAreDeviceCopyable
2447 : CheckFieldsAreDeviceCopyable<T, NumFieldsToCheck - 1> {
2448 using FieldT = decltype(__builtin_field_type(T, NumFieldsToCheck - 1));
2449 static_assert(is_device_copyable<FieldT>::value ||
2450 detail::IsDeprecatedDeviceCopyable<FieldT>::value,
2451 "The specified type is not device copyable");
2454 template <typename T> struct CheckFieldsAreDeviceCopyable<T, 0> {};
2458 template <typename T, unsigned NumBasesToCheck>
2459 struct CheckBasesAreDeviceCopyable
2460 : CheckBasesAreDeviceCopyable<T, NumBasesToCheck - 1> {
2461 using BaseT = decltype(__builtin_base_type(T, NumBasesToCheck - 1));
2462 static_assert(is_device_copyable<BaseT>::value ||
2463 detail::IsDeprecatedDeviceCopyable<BaseT>::value,
2464 "The specified type is not device copyable");
2467 template <typename T> struct CheckBasesAreDeviceCopyable<T, 0> {};
2480 template <typename FuncT>
2481 struct CheckDeviceCopyable
2482 : CheckFieldsAreDeviceCopyable<FuncT, __builtin_num_fields(FuncT)>,
2483 CheckBasesAreDeviceCopyable<FuncT, __builtin_num_bases(FuncT)> {};
2487 template <typename TransformedArgType, int Dims, typename KernelType>
2488 struct CheckDeviceCopyable<
2489 RoundedRangeKernel<TransformedArgType, Dims, KernelType>>
2490 : CheckDeviceCopyable<KernelType> {};
2492 template <typename TransformedArgType, int Dims, typename KernelType>
2493 struct CheckDeviceCopyable<
2494 RoundedRangeKernelWithKH<TransformedArgType, Dims, KernelType>>
2495 : CheckDeviceCopyable<KernelType> {};
2503 #undef __SYCL_ALIGNED_VAR