28 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H 29 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H 31 #if defined(__HCC__) && (__hcc_workweek__ < 16032) 32 #error("This version of HIP requires a newer version of HCC."); 37 #if !defined(_MSC_VER)|| __clang__ 38 #if __has_attribute(ext_vector_type) 39 #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) 41 #define __NATIVE_VECTOR__(n, ...) [n] 44 #if defined(__cplusplus) 47 #include <type_traits> 50 template<
typename,
typename,
unsigned int>
struct Scalar_accessor;
54 template<
typename T,
typename U,
unsigned int n>
55 struct is_integral<
hip_impl::Scalar_accessor<T, U, n>>
57 template<
typename T,
typename U,
unsigned int n>
58 struct is_floating_point<hip_impl::Scalar_accessor<T, U, n>>
59 : is_floating_point<T> {};
63 template<
typename T,
typename Vector,
unsigned int idx>
64 struct Scalar_accessor {
66 const Scalar_accessor* p;
69 operator const T*()
const noexcept {
70 return &
reinterpret_cast<const T*
>(p)[idx];
73 operator const T*()
const volatile noexcept {
74 return &
reinterpret_cast<const T*
>(p)[idx];
77 operator T*() noexcept {
78 return &
reinterpret_cast<T*
>(
79 const_cast<Scalar_accessor*
>(p))[idx];
82 operator T*()
volatile noexcept {
83 return &
reinterpret_cast<T*
>(
84 const_cast<Scalar_accessor*
>(p))[idx];
90 std::ostream& operator<<(std::ostream& os,
91 const Scalar_accessor& x) noexcept {
92 return os << x.data[idx];
96 std::istream& operator>>(std::istream& is,
97 Scalar_accessor& x) noexcept {
109 operator T() const noexcept {
return data[idx]; }
111 operator T() const volatile noexcept {
return data[idx]; }
113 #ifdef __HIP_ENABLE_VECTOR_SCALAR_ACCESSORY_ENUM_CONVERSION__ 118 typename std::enable_if<
119 !std::is_same<U, T>{} &&
122 T,
typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* =
nullptr>
124 operator U() const noexcept {
return static_cast<U
>(data[idx]); }
127 typename std::enable_if<
128 !std::is_same<U, T>{} &&
131 T,
typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* =
nullptr>
133 operator U() const volatile noexcept {
return static_cast<U
>(data[idx]); }
137 operator T&() noexcept {
138 return reinterpret_cast< 139 T (&)[sizeof(Vector) / sizeof(T)]
>(data)[idx];
142 operator volatile T&()
volatile noexcept {
143 return reinterpret_cast< 144 volatile T (&)[sizeof(Vector) / sizeof(T)]
>(data)[idx];
148 Address operator&() const noexcept {
return Address{
this}; }
151 Scalar_accessor& operator=(
const Scalar_accessor& x) noexcept {
152 data[idx] = x.data[idx];
157 Scalar_accessor& operator=(T x) noexcept {
163 volatile Scalar_accessor& operator=(T x)
volatile noexcept {
170 Scalar_accessor& operator++() noexcept {
175 T operator++(
int) noexcept {
181 Scalar_accessor& operator--() noexcept {
186 T operator--(
int) noexcept {
196 typename std::enable_if<
197 std::is_convertible<U, T>{}>::type* =
nullptr>
199 Scalar_accessor& operator+=(U x) noexcept {
205 typename std::enable_if<
206 std::is_convertible<U, T>{}>::type* =
nullptr>
208 Scalar_accessor& operator-=(U x) noexcept {
215 typename std::enable_if<
216 std::is_convertible<U, T>{}>::type* =
nullptr>
218 Scalar_accessor& operator*=(U x) noexcept {
224 typename std::enable_if<
225 std::is_convertible<U, T>{}>::type* =
nullptr>
227 Scalar_accessor& operator/=(U x) noexcept {
233 typename std::enable_if<std::is_convertible<U, T>{} &&
234 std::is_integral<U>{}>::type* =
nullptr>
236 Scalar_accessor& operator%=(U x) noexcept {
243 typename std::enable_if<std::is_convertible<U, T>{} &&
244 std::is_integral<U>{}>::type* =
nullptr>
246 Scalar_accessor& operator>>=(U x) noexcept {
252 typename std::enable_if<std::is_convertible<U, T>{} &&
253 std::is_integral<U>{}>::type* =
nullptr>
255 Scalar_accessor& operator<<=(U x) noexcept {
261 typename std::enable_if<std::is_convertible<U, T>{} &&
262 std::is_integral<U>{}>::type* =
nullptr>
264 Scalar_accessor& operator&=(U x) noexcept {
270 typename std::enable_if<std::is_convertible<U, T>{} &&
271 std::is_integral<U>{}>::type* =
nullptr>
273 Scalar_accessor& operator|=(U x) noexcept {
279 typename std::enable_if<std::is_convertible<U, T>{} &&
280 std::is_integral<U>{}>::type* =
nullptr>
282 Scalar_accessor& operator^=(U x) noexcept {
289 template<
typename T,
unsigned int n>
struct HIP_vector_base;
292 struct HIP_vector_base<T, 1> {
293 using Native_vec_ = T __NATIVE_VECTOR__(1, T);
297 #if __HIP_CLANG_ONLY__ 302 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
306 using value_type = T;
309 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
310 #if __has_attribute(ext_vector_type) 321 struct HIP_vector_base<T, 2> {
322 using Native_vec_ = T __NATIVE_VECTOR__(2, T);
326 #if __HIP_CLANG_ONLY__ 332 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
333 hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
337 using value_type = T;
340 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
341 #if __has_attribute(ext_vector_type) 353 struct HIP_vector_base<T, 3> {
359 Native_vec_() =
default;
363 Native_vec_(T x) noexcept : d{x, x, x} {}
366 Native_vec_(T x, T y, T z) noexcept : d{x, y, z} {}
369 Native_vec_(
const Native_vec_&) =
default;
372 Native_vec_(Native_vec_&&) =
default;
374 ~Native_vec_() =
default;
377 Native_vec_& operator=(
const Native_vec_&) =
default;
379 Native_vec_& operator=(Native_vec_&&) =
default;
382 T& operator[](
unsigned int idx) noexcept {
return d[idx]; }
384 T operator[](
unsigned int idx)
const noexcept {
return d[idx]; }
387 Native_vec_& operator+=(
const Native_vec_& x) noexcept
389 for (
auto i = 0u; i != 3u; ++i) d[i] += x.d[i];
393 Native_vec_& operator-=(
const Native_vec_& x) noexcept
395 for (
auto i = 0u; i != 3u; ++i) d[i] -= x.d[i];
400 Native_vec_& operator*=(
const Native_vec_& x) noexcept
402 for (
auto i = 0u; i != 3u; ++i) d[i] *= x.d[i];
406 Native_vec_& operator/=(
const Native_vec_& x) noexcept
408 for (
auto i = 0u; i != 3u; ++i) d[i] /= x.d[i];
414 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
416 Native_vec_ operator-() const noexcept
419 for (
auto&& x : r.d) x = -x;
425 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
427 Native_vec_ operator~() const noexcept
430 for (
auto&& x : r.d) x = ~x;
435 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
437 Native_vec_& operator%=(
const Native_vec_& x) noexcept
439 for (
auto i = 0u; i != 3u; ++i) d[i] %= x.d[i];
444 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
446 Native_vec_& operator^=(
const Native_vec_& x) noexcept
448 for (
auto i = 0u; i != 3u; ++i) d[i] ^= x.d[i];
453 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
455 Native_vec_& operator|=(
const Native_vec_& x) noexcept
457 for (
auto i = 0u; i != 3u; ++i) d[i] |= x.d[i];
462 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
464 Native_vec_& operator&=(
const Native_vec_& x) noexcept
466 for (
auto i = 0u; i != 3u; ++i) d[i] &= x.d[i];
471 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
473 Native_vec_& operator>>=(
const Native_vec_& x) noexcept
475 for (
auto i = 0u; i != 3u; ++i) d[i] >>= x.d[i];
480 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
482 Native_vec_& operator<<=(
const Native_vec_& x) noexcept
484 for (
auto i = 0u; i != 3u; ++i) d[i] <<= x.d[i];
488 using Vec3_cmp =
int __attribute__((vector_size(4 *
sizeof(
int))));
490 Vec3_cmp operator==(
const Native_vec_& x)
const noexcept
492 return Vec3_cmp{d[0] == x.d[0], d[1] == x.d[1], d[2] == x.d[2]};
505 using value_type = T;
509 struct HIP_vector_base<T, 4> {
510 using Native_vec_ = T __NATIVE_VECTOR__(4, T);
514 #if __HIP_CLANG_ONLY__ 522 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
523 hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
524 hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
525 hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
529 using value_type = T;
532 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
533 #if __has_attribute(ext_vector_type) 546 template<
typename T,
unsigned int rank>
547 struct HIP_vector_type :
public HIP_vector_base<T, rank> {
548 using HIP_vector_base<T, rank>::data;
549 using typename HIP_vector_base<T, rank>::Native_vec_;
552 HIP_vector_type() =
default;
555 typename std::enable_if<
556 std::is_convertible<U, T>{}>::type* =
nullptr>
558 HIP_vector_type(U x) noexcept
560 for (
auto i = 0u; i != rank; ++i) data[i] = x;
564 typename std::enable_if<
565 (rank > 1) &&
sizeof...(Us) == rank>::type* =
nullptr>
567 HIP_vector_type(Us... xs) noexcept
569 #if __has_attribute(ext_vector_type) 570 new (&data) Native_vec_{
static_cast<T
>(xs)...};
572 new (&data) std::array<T, rank>{
static_cast<T
>(xs)...};
576 HIP_vector_type(
const HIP_vector_type&) =
default;
578 HIP_vector_type(HIP_vector_type&&) =
default;
580 ~HIP_vector_type() =
default;
583 HIP_vector_type& operator=(
const HIP_vector_type&) =
default;
585 HIP_vector_type& operator=(HIP_vector_type&&) =
default;
589 HIP_vector_type& operator++() noexcept
591 return *
this += HIP_vector_type{1};
594 HIP_vector_type operator++(
int) noexcept
602 HIP_vector_type& operator--() noexcept
604 return *
this -= HIP_vector_type{1};
607 HIP_vector_type operator--(
int) noexcept
615 HIP_vector_type& operator+=(
const HIP_vector_type& x) noexcept
622 typename std::enable_if<
623 std::is_convertible<U, T>{}>::type* =
nullptr>
625 HIP_vector_type& operator+=(U x) noexcept
627 return *
this += HIP_vector_type{x};
631 HIP_vector_type& operator-=(
const HIP_vector_type& x) noexcept
638 typename std::enable_if<
639 std::is_convertible<U, T>{}>::type* =
nullptr>
641 HIP_vector_type& operator-=(U x) noexcept
643 return *
this -= HIP_vector_type{x};
647 HIP_vector_type& operator*=(
const HIP_vector_type& x) noexcept
654 typename std::enable_if<
655 std::is_convertible<U, T>{}>::type* =
nullptr>
657 HIP_vector_type& operator*=(U x) noexcept
659 return *
this *= HIP_vector_type{x};
663 HIP_vector_type& operator/=(
const HIP_vector_type& x) noexcept
670 typename std::enable_if<
671 std::is_convertible<U, T>{}>::type* =
nullptr>
673 HIP_vector_type& operator/=(U x) noexcept
675 return *
this /= HIP_vector_type{x};
680 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
682 HIP_vector_type operator-() noexcept
685 tmp.data = -tmp.data;
691 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
693 HIP_vector_type operator~() noexcept
695 HIP_vector_type r{*
this};
702 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
704 HIP_vector_type& operator%=(
const HIP_vector_type& x) noexcept
712 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
714 HIP_vector_type& operator^=(
const HIP_vector_type& x) noexcept
722 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
724 HIP_vector_type& operator|=(
const HIP_vector_type& x) noexcept
732 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
734 HIP_vector_type& operator&=(
const HIP_vector_type& x) noexcept
742 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
744 HIP_vector_type& operator>>=(
const HIP_vector_type& x) noexcept
752 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
754 HIP_vector_type& operator<<=(
const HIP_vector_type& x) noexcept
761 template<
typename T,
unsigned int n>
763 HIP_vector_type<T, n> operator+(
764 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
766 return HIP_vector_type<T, n>{x} += y;
768 template<
typename T,
unsigned int n,
typename U>
770 HIP_vector_type<T, n> operator+(
771 const HIP_vector_type<T, n>& x, U y) noexcept
773 return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
775 template<
typename T,
unsigned int n,
typename U>
777 HIP_vector_type<T, n> operator+(
778 U x,
const HIP_vector_type<T, n>& y) noexcept
780 return HIP_vector_type<T, n>{x} += y;
783 template<
typename T,
unsigned int n>
785 HIP_vector_type<T, n> operator-(
786 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
788 return HIP_vector_type<T, n>{x} -= y;
790 template<
typename T,
unsigned int n,
typename U>
792 HIP_vector_type<T, n> operator-(
793 const HIP_vector_type<T, n>& x, U y) noexcept
795 return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
797 template<
typename T,
unsigned int n,
typename U>
799 HIP_vector_type<T, n> operator-(
800 U x,
const HIP_vector_type<T, n>& y) noexcept
802 return HIP_vector_type<T, n>{x} -= y;
805 template<
typename T,
unsigned int n>
807 HIP_vector_type<T, n> operator*(
808 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
810 return HIP_vector_type<T, n>{x} *= y;
812 template<
typename T,
unsigned int n,
typename U>
814 HIP_vector_type<T, n> operator*(
815 const HIP_vector_type<T, n>& x, U y) noexcept
817 return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
819 template<
typename T,
unsigned int n,
typename U>
821 HIP_vector_type<T, n> operator*(
822 U x,
const HIP_vector_type<T, n>& y) noexcept
824 return HIP_vector_type<T, n>{x} *= y;
827 template<
typename T,
unsigned int n>
829 HIP_vector_type<T, n> operator/(
830 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
832 return HIP_vector_type<T, n>{x} /= y;
834 template<
typename T,
unsigned int n,
typename U>
836 HIP_vector_type<T, n> operator/(
837 const HIP_vector_type<T, n>& x, U y) noexcept
839 return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
841 template<
typename T,
unsigned int n,
typename U>
843 HIP_vector_type<T, n> operator/(
844 U x,
const HIP_vector_type<T, n>& y) noexcept
846 return HIP_vector_type<T, n>{x} /= y;
849 template<
typename T,
unsigned int n>
852 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
854 auto tmp = x.data == y.data;
855 for (
auto i = 0u; i != n; ++i)
if (tmp[i] == 0)
return false;
858 template<
typename T,
unsigned int n,
typename U>
860 bool operator==(
const HIP_vector_type<T, n>& x, U y) noexcept
862 return x == HIP_vector_type<T, n>{y};
864 template<
typename T,
unsigned int n,
typename U>
866 bool operator==(U x,
const HIP_vector_type<T, n>& y) noexcept
868 return HIP_vector_type<T, n>{x} == y;
871 template<
typename T,
unsigned int n>
874 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
878 template<
typename T,
unsigned int n,
typename U>
880 bool operator!=(
const HIP_vector_type<T, n>& x, U y) noexcept
884 template<
typename T,
unsigned int n,
typename U>
886 bool operator!=(U x,
const HIP_vector_type<T, n>& y) noexcept
894 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
896 HIP_vector_type<T, n> operator%(
897 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
899 return HIP_vector_type<T, n>{x} %= y;
905 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
907 HIP_vector_type<T, n> operator%(
908 const HIP_vector_type<T, n>& x, U y) noexcept
910 return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
916 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
918 HIP_vector_type<T, n> operator%(
919 U x,
const HIP_vector_type<T, n>& y) noexcept
921 return HIP_vector_type<T, n>{x} %= y;
927 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
929 HIP_vector_type<T, n> operator^(
930 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
932 return HIP_vector_type<T, n>{x} ^= y;
938 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
940 HIP_vector_type<T, n> operator^(
941 const HIP_vector_type<T, n>& x, U y) noexcept
943 return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
949 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
951 HIP_vector_type<T, n> operator^(
952 U x,
const HIP_vector_type<T, n>& y) noexcept
954 return HIP_vector_type<T, n>{x} ^= y;
960 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
962 HIP_vector_type<T, n> operator|(
963 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
965 return HIP_vector_type<T, n>{x} |= y;
971 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
973 HIP_vector_type<T, n> operator|(
974 const HIP_vector_type<T, n>& x, U y) noexcept
976 return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
982 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
984 HIP_vector_type<T, n> operator|(
985 U x,
const HIP_vector_type<T, n>& y) noexcept
987 return HIP_vector_type<T, n>{x} |= y;
993 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
995 HIP_vector_type<T, n> operator&(
996 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
998 return HIP_vector_type<T, n>{x} &= y;
1004 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1006 HIP_vector_type<T, n> operator&(
1007 const HIP_vector_type<T, n>& x, U y) noexcept
1009 return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
1015 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1017 HIP_vector_type<T, n> operator&(
1018 U x,
const HIP_vector_type<T, n>& y) noexcept
1020 return HIP_vector_type<T, n>{x} &= y;
1026 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1028 HIP_vector_type<T, n> operator>>(
1029 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1031 return HIP_vector_type<T, n>{x} >>= y;
1037 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1039 HIP_vector_type<T, n> operator>>(
1040 const HIP_vector_type<T, n>& x, U y) noexcept
1042 return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
1048 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1050 HIP_vector_type<T, n> operator>>(
1051 U x,
const HIP_vector_type<T, n>& y) noexcept
1053 return HIP_vector_type<T, n>{x} >>= y;
1059 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1061 HIP_vector_type<T, n> operator<<(
1062 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1064 return HIP_vector_type<T, n>{x} <<= y;
1070 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1072 HIP_vector_type<T, n> operator<<(
1073 const HIP_vector_type<T, n>& x, U y) noexcept
1075 return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
1081 typename std::enable_if<std::is_arithmetic<U>::value>::type,
1082 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1084 HIP_vector_type<T, n> operator<<(
1085 U x,
const HIP_vector_type<T, n>& y) noexcept
1087 return HIP_vector_type<T, n>{x} <<= y;
1090 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 1091 using CUDA_name##1 = HIP_vector_type<T, 1>;\ 1092 using CUDA_name##2 = HIP_vector_type<T, 2>;\ 1093 using CUDA_name##3 = HIP_vector_type<T, 3>;\ 1094 using CUDA_name##4 = HIP_vector_type<T, 4>; 1096 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 1117 __MAKE_VECTOR_TYPE__(uchar,
unsigned char);
1118 __MAKE_VECTOR_TYPE__(
char,
char);
1119 __MAKE_VECTOR_TYPE__(ushort,
unsigned short);
1120 __MAKE_VECTOR_TYPE__(
short,
short);
1121 __MAKE_VECTOR_TYPE__(uint,
unsigned int);
1122 __MAKE_VECTOR_TYPE__(
int,
int);
1123 __MAKE_VECTOR_TYPE__(ulong,
unsigned long);
1124 __MAKE_VECTOR_TYPE__(
long,
long);
1125 __MAKE_VECTOR_TYPE__(ulonglong,
unsigned long long);
1126 __MAKE_VECTOR_TYPE__(longlong,
long long);
1127 __MAKE_VECTOR_TYPE__(
float,
float);
1128 __MAKE_VECTOR_TYPE__(
double,
double);
1131 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 1132 static inline __device__ __host__ \ 1133 type make_##type(comp x) { type r{x}; return r; } 1135 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 1136 static inline __device__ __host__ \ 1137 type make_##type(comp x, comp y) { type r{x, y}; return r; } 1139 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 1140 static inline __device__ __host__ \ 1141 type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; } 1143 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 1144 static inline __device__ __host__ \ 1145 type make_##type(comp x, comp y, comp z, comp w) { \ 1146 type r{x, y, z, w}; \ 1150 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 1151 static inline __device__ __host__ \ 1152 type make_##type(comp x) { type r; r.x =x; return r; } 1154 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 1155 static inline __device__ __host__ \ 1156 type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; } 1158 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 1159 static inline __device__ __host__ \ 1160 type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; } 1162 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 1163 static inline __device__ __host__ \ 1164 type make_##type(comp x, comp y, comp z, comp w) { \ 1165 type r; r.x=x; r.y=y; r.z=z; r.w=w; \ 1170 DECLOP_MAKE_ONE_COMPONENT(
unsigned char, uchar1);
1171 DECLOP_MAKE_TWO_COMPONENT(
unsigned char, uchar2);
1172 DECLOP_MAKE_THREE_COMPONENT(
unsigned char, uchar3);
1173 DECLOP_MAKE_FOUR_COMPONENT(
unsigned char, uchar4);
1175 DECLOP_MAKE_ONE_COMPONENT(
signed char, char1);
1176 DECLOP_MAKE_TWO_COMPONENT(
signed char, char2);
1177 DECLOP_MAKE_THREE_COMPONENT(
signed char, char3);
1178 DECLOP_MAKE_FOUR_COMPONENT(
signed char, char4);
1180 DECLOP_MAKE_ONE_COMPONENT(
unsigned short, ushort1);
1181 DECLOP_MAKE_TWO_COMPONENT(
unsigned short, ushort2);
1182 DECLOP_MAKE_THREE_COMPONENT(
unsigned short, ushort3);
1183 DECLOP_MAKE_FOUR_COMPONENT(
unsigned short, ushort4);
1185 DECLOP_MAKE_ONE_COMPONENT(
signed short, short1);
1186 DECLOP_MAKE_TWO_COMPONENT(
signed short, short2);
1187 DECLOP_MAKE_THREE_COMPONENT(
signed short, short3);
1188 DECLOP_MAKE_FOUR_COMPONENT(
signed short, short4);
1190 DECLOP_MAKE_ONE_COMPONENT(
unsigned int, uint1);
1191 DECLOP_MAKE_TWO_COMPONENT(
unsigned int, uint2);
1192 DECLOP_MAKE_THREE_COMPONENT(
unsigned int, uint3);
1193 DECLOP_MAKE_FOUR_COMPONENT(
unsigned int, uint4);
1195 DECLOP_MAKE_ONE_COMPONENT(
signed int, int1);
1196 DECLOP_MAKE_TWO_COMPONENT(
signed int, int2);
1197 DECLOP_MAKE_THREE_COMPONENT(
signed int, int3);
1198 DECLOP_MAKE_FOUR_COMPONENT(
signed int, int4);
1200 DECLOP_MAKE_ONE_COMPONENT(
float, float1);
1201 DECLOP_MAKE_TWO_COMPONENT(
float, float2);
1202 DECLOP_MAKE_THREE_COMPONENT(
float, float3);
1203 DECLOP_MAKE_FOUR_COMPONENT(
float, float4);
1205 DECLOP_MAKE_ONE_COMPONENT(
double, double1);
1206 DECLOP_MAKE_TWO_COMPONENT(
double, double2);
1207 DECLOP_MAKE_THREE_COMPONENT(
double, double3);
1208 DECLOP_MAKE_FOUR_COMPONENT(
double, double4);
1210 DECLOP_MAKE_ONE_COMPONENT(
unsigned long, ulong1);
1211 DECLOP_MAKE_TWO_COMPONENT(
unsigned long, ulong2);
1212 DECLOP_MAKE_THREE_COMPONENT(
unsigned long, ulong3);
1213 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long, ulong4);
1215 DECLOP_MAKE_ONE_COMPONENT(
signed long, long1);
1216 DECLOP_MAKE_TWO_COMPONENT(
signed long, long2);
1217 DECLOP_MAKE_THREE_COMPONENT(
signed long, long3);
1218 DECLOP_MAKE_FOUR_COMPONENT(
signed long, long4);
1220 DECLOP_MAKE_ONE_COMPONENT(
unsigned long long, ulonglong1);
1221 DECLOP_MAKE_TWO_COMPONENT(
unsigned long long, ulonglong2);
1222 DECLOP_MAKE_THREE_COMPONENT(
unsigned long long, ulonglong3);
1223 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long long, ulonglong4);
1225 DECLOP_MAKE_ONE_COMPONENT(
signed long long, longlong1);
1226 DECLOP_MAKE_TWO_COMPONENT(
signed long long, longlong2);
1227 DECLOP_MAKE_THREE_COMPONENT(
signed long long, longlong3);
1228 DECLOP_MAKE_FOUR_COMPONENT(
signed long long, longlong4);
1229 #else // defined(_MSC_VER) 1230 #include <mmintrin.h> 1231 #include <xmmintrin.h> 1232 #include <emmintrin.h> 1233 #include <immintrin.h> 1235 typedef union {
char data; } char1;
1236 typedef union {
char data[2]; } char2;
1237 typedef union {
char data[4]; } char4;
1238 typedef union { char4 data; } char3;
1239 typedef union { __m64 data; } char8;
1240 typedef union { __m128i data; } char16;
1242 typedef union {
unsigned char data; } uchar1;
1243 typedef union {
unsigned char data[2]; } uchar2;
1244 typedef union {
unsigned char data[4]; } uchar4;
1245 typedef union { uchar4 data; } uchar3;
1246 typedef union { __m64 data; } uchar8;
1247 typedef union { __m128i data; } uchar16;
1249 typedef union {
short data; } short1;
1250 typedef union {
short data[2]; } short2;
1251 typedef union { __m64 data; } short4;
1252 typedef union { short4 data; } short3;
1253 typedef union { __m128i data; } short8;
1254 typedef union { __m128i data[2]; } short16;
1256 typedef union {
unsigned short data; } ushort1;
1257 typedef union {
unsigned short data[2]; } ushort2;
1258 typedef union { __m64 data; } ushort4;
1259 typedef union { ushort4 data; } ushort3;
1260 typedef union { __m128i data; } ushort8;
1261 typedef union { __m128i data[2]; } ushort16;
1263 typedef union {
int data; } int1;
1264 typedef union { __m64 data; } int2;
1265 typedef union { __m128i data; } int4;
1266 typedef union { int4 data; } int3;
1267 typedef union { __m128i data[2]; } int8;
1268 typedef union { __m128i data[4];} int16;
1270 typedef union {
unsigned int data; } uint1;
1271 typedef union { __m64 data; } uint2;
1272 typedef union { __m128i data; } uint4;
1273 typedef union { uint4 data; } uint3;
1274 typedef union { __m128i data[2]; } uint8;
1275 typedef union { __m128i data[4]; } uint16;
1277 #if !defined(_WIN64) 1278 typedef union {
int data; } long1;
1279 typedef union { __m64 data; } long2;
1280 typedef union { __m128i data; } long4;
1281 typedef union { long4 data; } long3;
1282 typedef union { __m128i data[2]; } long8;
1283 typedef union { __m128i data[4]; } long16;
1285 typedef union {
unsigned int data; } ulong1;
1286 typedef union { __m64 data; } ulong2;
1287 typedef union { __m128i data; } ulong4;
1288 typedef union { ulong4 data; } ulong3;
1289 typedef union { __m128i data[2]; } ulong8;
1290 typedef union { __m128i data[4]; } ulong16;
1291 #else // defined(_WIN64) 1292 typedef union { __m64 data; } long1;
1293 typedef union { __m128i data; } long2;
1294 typedef union { __m128i data[2]; } long4;
1295 typedef union { long4 data; } long3;
1296 typedef union { __m128i data[4]; } long8;
1297 typedef union { __m128i data[8]; } long16;
1299 typedef union { __m64 data; } ulong1;
1300 typedef union { __m128i data; } ulong2;
1301 typedef union { __m128i data[2]; } ulong4;
1302 typedef union { ulong4 data; } ulong3;
1303 typedef union { __m128i data[4]; } ulong8;
1304 typedef union { __m128i data[8]; } ulong16;
1305 #endif // defined(_WIN64) 1307 typedef union { __m64 data; } longlong1;
1308 typedef union { __m128i data; } longlong2;
1309 typedef union { __m128i data[2]; } longlong4;
1310 typedef union { longlong4 data; } longlong3;
1311 typedef union { __m128i data[4]; } longlong8;
1312 typedef union { __m128i data[8]; } longlong16;
1314 typedef union { __m64 data; } ulonglong1;
1315 typedef union { __m128i data; } ulonglong2;
1316 typedef union { __m128i data[2]; } ulonglong4;
1317 typedef union { ulonglong4 data; } ulonglong3;
1318 typedef union { __m128i data[4]; } ulonglong8;
1319 typedef union { __m128i data[8]; } ulonglong16;
1321 typedef union {
float data; } float1;
1322 typedef union { __m64 data; } float2;
1323 typedef union { __m128 data; } float4;
1324 typedef union { float4 data; } float3;
1325 typedef union { __m256 data; } float8;
1326 typedef union { __m256 data[2]; } float16;
1328 typedef union {
double data; } double1;
1329 typedef union { __m128d data; } double2;
1330 typedef union { __m256d data; } double4;
1331 typedef union { double4 data; } double3;
1332 typedef union { __m256d data[2]; } double8;
1333 typedef union { __m256d data[4]; } double16;
1335 #endif // defined(_MSC_VER)
#define __host__
Definition: host_defines.h:41
Definition: hip_runtime.h:202