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, T) T __attribute__((ext_vector_type(n))) 41 #define __NATIVE_VECTOR__(n, T) T[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 {
290 unsigned int next_pot(
unsigned int x) {
292 return 1u << (32u - __builtin_clz(x - 1u));
296 template<
typename T,
unsigned int n>
struct HIP_vector_base;
299 struct HIP_vector_base<T, 1> {
300 using Native_vec_ = __NATIVE_VECTOR__(1, T);
304 #if __HIP_CLANG_ONLY__ 309 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
313 using value_type = T;
316 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
317 #if __has_attribute(ext_vector_type) 328 struct HIP_vector_base<T, 2> {
329 using Native_vec_ = __NATIVE_VECTOR__(2, T);
332 #if !__has_attribute(ext_vector_type) 333 alignas(hip_impl::next_pot(2 *
sizeof(T)))
337 #if __HIP_CLANG_ONLY__ 343 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
344 hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
348 using value_type = T;
351 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
352 #if __has_attribute(ext_vector_type) 364 struct HIP_vector_base<T, 3> {
370 Native_vec_() =
default;
374 Native_vec_(T x) noexcept : d{x, x, x} {}
377 Native_vec_(T x, T y, T z) noexcept : d{x, y, z} {}
380 Native_vec_(
const Native_vec_&) =
default;
383 Native_vec_(Native_vec_&&) =
default;
385 ~Native_vec_() =
default;
388 Native_vec_& operator=(
const Native_vec_&) =
default;
390 Native_vec_& operator=(Native_vec_&&) =
default;
393 T& operator[](
unsigned int idx) noexcept {
return d[idx]; }
395 T operator[](
unsigned int idx)
const noexcept {
return d[idx]; }
398 Native_vec_& operator+=(
const Native_vec_& x) noexcept
400 for (
auto i = 0u; i != 3u; ++i) d[i] += x.d[i];
404 Native_vec_& operator-=(
const Native_vec_& x) noexcept
406 for (
auto i = 0u; i != 3u; ++i) d[i] -= x.d[i];
411 Native_vec_& operator*=(
const Native_vec_& x) noexcept
413 for (
auto i = 0u; i != 3u; ++i) d[i] *= x.d[i];
417 Native_vec_& operator/=(
const Native_vec_& x) noexcept
419 for (
auto i = 0u; i != 3u; ++i) d[i] /= x.d[i];
425 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
427 Native_vec_ operator-() const noexcept
430 for (
auto&& x : r.d) x = -x;
436 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
438 Native_vec_ operator~() const noexcept
441 for (
auto&& x : r.d) x = ~x;
446 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
448 Native_vec_& operator%=(
const Native_vec_& x) noexcept
450 for (
auto i = 0u; i != 3u; ++i) d[i] %= x.d[i];
455 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
457 Native_vec_& operator^=(
const Native_vec_& x) noexcept
459 for (
auto i = 0u; i != 3u; ++i) d[i] ^= x.d[i];
464 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
466 Native_vec_& operator|=(
const Native_vec_& x) noexcept
468 for (
auto i = 0u; i != 3u; ++i) d[i] |= x.d[i];
473 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
475 Native_vec_& operator&=(
const Native_vec_& x) noexcept
477 for (
auto i = 0u; i != 3u; ++i) d[i] &= x.d[i];
482 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
484 Native_vec_& operator>>=(
const Native_vec_& x) noexcept
486 for (
auto i = 0u; i != 3u; ++i) d[i] >>= x.d[i];
491 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
493 Native_vec_& operator<<=(
const Native_vec_& x) noexcept
495 for (
auto i = 0u; i != 3u; ++i) d[i] <<= x.d[i];
499 using Vec3_cmp =
int __attribute__((vector_size(4 *
sizeof(
int))));
501 Vec3_cmp operator==(
const Native_vec_& x)
const noexcept
503 return Vec3_cmp{d[0] == x.d[0], d[1] == x.d[1], d[2] == x.d[2]};
516 using value_type = T;
520 struct HIP_vector_base<T, 4> {
521 using Native_vec_ = __NATIVE_VECTOR__(4, T);
524 #if !__has_attribute(ext_vector_type) 525 alignas(hip_impl::next_pot(4 *
sizeof(T)))
529 #if __HIP_CLANG_ONLY__ 537 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
538 hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
539 hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
540 hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
544 using value_type = T;
547 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
548 #if __has_attribute(ext_vector_type) 561 template<
typename T,
unsigned int rank>
562 struct HIP_vector_type :
public HIP_vector_base<T, rank> {
563 using HIP_vector_base<T, rank>::data;
564 using typename HIP_vector_base<T, rank>::Native_vec_;
567 HIP_vector_type() =
default;
570 typename std::enable_if<
571 std::is_convertible<U, T>{}>::type* =
nullptr>
573 HIP_vector_type(U x) noexcept
575 for (
auto i = 0u; i != rank; ++i) data[i] = x;
579 typename std::enable_if<
580 (rank > 1) &&
sizeof...(Us) == rank>::type* =
nullptr>
582 HIP_vector_type(Us... xs) noexcept
584 #if __has_attribute(ext_vector_type) 585 new (&data) Native_vec_{
static_cast<T
>(xs)...};
587 new (&data) std::array<T, rank>{
static_cast<T
>(xs)...};
591 HIP_vector_type(
const HIP_vector_type&) =
default;
593 HIP_vector_type(HIP_vector_type&&) =
default;
595 ~HIP_vector_type() =
default;
598 HIP_vector_type& operator=(
const HIP_vector_type&) =
default;
600 HIP_vector_type& operator=(HIP_vector_type&&) =
default;
604 HIP_vector_type& operator++() noexcept
606 return *
this += HIP_vector_type{1};
609 HIP_vector_type operator++(
int) noexcept
617 HIP_vector_type& operator--() noexcept
619 return *
this -= HIP_vector_type{1};
622 HIP_vector_type operator--(
int) noexcept
630 HIP_vector_type& operator+=(
const HIP_vector_type& x) noexcept
637 typename std::enable_if<
638 std::is_convertible<U, T>{}>::type* =
nullptr>
640 HIP_vector_type& operator+=(U x) noexcept
642 return *
this += HIP_vector_type{x};
646 HIP_vector_type& operator-=(
const HIP_vector_type& x) noexcept
653 typename std::enable_if<
654 std::is_convertible<U, T>{}>::type* =
nullptr>
656 HIP_vector_type& operator-=(U x) noexcept
658 return *
this -= HIP_vector_type{x};
662 HIP_vector_type& operator*=(
const HIP_vector_type& x) noexcept
669 typename std::enable_if<
670 std::is_convertible<U, T>{}>::type* =
nullptr>
672 HIP_vector_type& operator*=(U x) noexcept
674 return *
this *= HIP_vector_type{x};
678 HIP_vector_type& operator/=(
const HIP_vector_type& x) noexcept
685 typename std::enable_if<
686 std::is_convertible<U, T>{}>::type* =
nullptr>
688 HIP_vector_type& operator/=(U x) noexcept
690 return *
this /= HIP_vector_type{x};
695 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
697 HIP_vector_type operator-() noexcept
700 tmp.data = -tmp.data;
706 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
708 HIP_vector_type operator~() noexcept
710 HIP_vector_type r{*
this};
717 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
719 HIP_vector_type& operator%=(
const HIP_vector_type& x) noexcept
727 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
729 HIP_vector_type& operator^=(
const HIP_vector_type& x) noexcept
737 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
739 HIP_vector_type& operator|=(
const HIP_vector_type& x) noexcept
747 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
749 HIP_vector_type& operator&=(
const HIP_vector_type& x) noexcept
757 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
759 HIP_vector_type& operator>>=(
const HIP_vector_type& x) noexcept
767 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
769 HIP_vector_type& operator<<=(
const HIP_vector_type& x) noexcept
776 template<
typename T,
unsigned int n>
778 HIP_vector_type<T, n> operator+(
779 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
781 return HIP_vector_type<T, n>{x} += y;
783 template<
typename T,
unsigned int n,
typename U>
785 HIP_vector_type<T, n> operator+(
786 const HIP_vector_type<T, n>& x, U y) noexcept
788 return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
790 template<
typename T,
unsigned int n,
typename U>
792 HIP_vector_type<T, n> operator+(
793 U x,
const HIP_vector_type<T, n>& y) noexcept
795 return HIP_vector_type<T, n>{x} += y;
798 template<
typename T,
unsigned int n>
800 HIP_vector_type<T, n> operator-(
801 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
803 return HIP_vector_type<T, n>{x} -= y;
805 template<
typename T,
unsigned int n,
typename U>
807 HIP_vector_type<T, n> operator-(
808 const HIP_vector_type<T, n>& x, U y) noexcept
810 return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
812 template<
typename T,
unsigned int n,
typename U>
814 HIP_vector_type<T, n> operator-(
815 U x,
const HIP_vector_type<T, n>& y) noexcept
817 return HIP_vector_type<T, n>{x} -= y;
820 template<
typename T,
unsigned int n>
822 HIP_vector_type<T, n> operator*(
823 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
825 return HIP_vector_type<T, n>{x} *= y;
827 template<
typename T,
unsigned int n,
typename U>
829 HIP_vector_type<T, n> operator*(
830 const HIP_vector_type<T, n>& x, U y) noexcept
832 return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
834 template<
typename T,
unsigned int n,
typename U>
836 HIP_vector_type<T, n> operator*(
837 U x,
const HIP_vector_type<T, n>& y) noexcept
839 return HIP_vector_type<T, n>{x} *= y;
842 template<
typename T,
unsigned int n>
844 HIP_vector_type<T, n> operator/(
845 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
847 return HIP_vector_type<T, n>{x} /= y;
849 template<
typename T,
unsigned int n,
typename U>
851 HIP_vector_type<T, n> operator/(
852 const HIP_vector_type<T, n>& x, U y) noexcept
854 return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
856 template<
typename T,
unsigned int n,
typename U>
858 HIP_vector_type<T, n> operator/(
859 U x,
const HIP_vector_type<T, n>& y) noexcept
861 return HIP_vector_type<T, n>{x} /= y;
864 template<
typename T,
unsigned int n>
867 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
869 auto tmp = x.data == y.data;
870 for (
auto i = 0u; i != n; ++i)
if (tmp[i] == 0)
return false;
873 template<
typename T,
unsigned int n,
typename U>
875 bool operator==(
const HIP_vector_type<T, n>& x, U y) noexcept
877 return x == HIP_vector_type<T, n>{y};
879 template<
typename T,
unsigned int n,
typename U>
881 bool operator==(U x,
const HIP_vector_type<T, n>& y) noexcept
883 return HIP_vector_type<T, n>{x} == y;
886 template<
typename T,
unsigned int n>
889 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
893 template<
typename T,
unsigned int n,
typename U>
895 bool operator!=(
const HIP_vector_type<T, n>& x, U y) noexcept
899 template<
typename T,
unsigned int n,
typename U>
901 bool operator!=(U x,
const HIP_vector_type<T, n>& y) noexcept
909 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
911 HIP_vector_type<T, n> operator%(
912 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
914 return HIP_vector_type<T, n>{x} %= y;
920 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
922 HIP_vector_type<T, n> operator%(
923 const HIP_vector_type<T, n>& x, U y) noexcept
925 return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
931 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
933 HIP_vector_type<T, n> operator%(
934 U x,
const HIP_vector_type<T, n>& y) noexcept
936 return HIP_vector_type<T, n>{x} %= y;
942 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
944 HIP_vector_type<T, n> operator^(
945 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
947 return HIP_vector_type<T, n>{x} ^= y;
953 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
955 HIP_vector_type<T, n> operator^(
956 const HIP_vector_type<T, n>& x, U y) noexcept
958 return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
964 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
966 HIP_vector_type<T, n> operator^(
967 U x,
const HIP_vector_type<T, n>& y) noexcept
969 return HIP_vector_type<T, n>{x} ^= y;
975 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
977 HIP_vector_type<T, n> operator|(
978 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
980 return HIP_vector_type<T, n>{x} |= y;
986 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
988 HIP_vector_type<T, n> operator|(
989 const HIP_vector_type<T, n>& x, U y) noexcept
991 return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
997 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
999 HIP_vector_type<T, n> operator|(
1000 U x,
const HIP_vector_type<T, n>& y) noexcept
1002 return HIP_vector_type<T, n>{x} |= y;
1008 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1010 HIP_vector_type<T, n> operator&(
1011 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1013 return HIP_vector_type<T, n>{x} &= y;
1019 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1021 HIP_vector_type<T, n> operator&(
1022 const HIP_vector_type<T, n>& x, U y) noexcept
1024 return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
1030 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1032 HIP_vector_type<T, n> operator&(
1033 U x,
const HIP_vector_type<T, n>& y) noexcept
1035 return HIP_vector_type<T, n>{x} &= y;
1041 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1043 HIP_vector_type<T, n> operator>>(
1044 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1046 return HIP_vector_type<T, n>{x} >>= y;
1052 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1054 HIP_vector_type<T, n> operator>>(
1055 const HIP_vector_type<T, n>& x, U y) noexcept
1057 return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
1063 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1065 HIP_vector_type<T, n> operator>>(
1066 U x,
const HIP_vector_type<T, n>& y) noexcept
1068 return HIP_vector_type<T, n>{x} >>= y;
1074 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1076 HIP_vector_type<T, n> operator<<(
1077 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1079 return HIP_vector_type<T, n>{x} <<= y;
1085 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1087 HIP_vector_type<T, n> operator<<(
1088 const HIP_vector_type<T, n>& x, U y) noexcept
1090 return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
1096 typename std::enable_if<std::is_arithmetic<U>::value>::type,
1097 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1099 HIP_vector_type<T, n> operator<<(
1100 U x,
const HIP_vector_type<T, n>& y) noexcept
1102 return HIP_vector_type<T, n>{x} <<= y;
1105 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 1106 using CUDA_name##1 = HIP_vector_type<T, 1>;\ 1107 using CUDA_name##2 = HIP_vector_type<T, 2>;\ 1108 using CUDA_name##3 = HIP_vector_type<T, 3>;\ 1109 using CUDA_name##4 = HIP_vector_type<T, 4>; 1111 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 1132 __MAKE_VECTOR_TYPE__(uchar,
unsigned char);
1133 __MAKE_VECTOR_TYPE__(
char,
char);
1134 __MAKE_VECTOR_TYPE__(ushort,
unsigned short);
1135 __MAKE_VECTOR_TYPE__(
short,
short);
1136 __MAKE_VECTOR_TYPE__(uint,
unsigned int);
1137 __MAKE_VECTOR_TYPE__(
int,
int);
1138 __MAKE_VECTOR_TYPE__(ulong,
unsigned long);
1139 __MAKE_VECTOR_TYPE__(
long,
long);
1140 __MAKE_VECTOR_TYPE__(ulonglong,
unsigned long long);
1141 __MAKE_VECTOR_TYPE__(longlong,
long long);
1142 __MAKE_VECTOR_TYPE__(
float,
float);
1143 __MAKE_VECTOR_TYPE__(
double,
double);
1146 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 1147 static inline __device__ __host__ \ 1148 type make_##type(comp x) { type r{x}; return r; } 1150 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 1151 static inline __device__ __host__ \ 1152 type make_##type(comp x, comp y) { type r{x, y}; return r; } 1154 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 1155 static inline __device__ __host__ \ 1156 type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; } 1158 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 1159 static inline __device__ __host__ \ 1160 type make_##type(comp x, comp y, comp z, comp w) { \ 1161 type r{x, y, z, w}; \ 1165 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 1166 static inline __device__ __host__ \ 1167 type make_##type(comp x) { type r; r.x =x; return r; } 1169 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 1170 static inline __device__ __host__ \ 1171 type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; } 1173 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 1174 static inline __device__ __host__ \ 1175 type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; } 1177 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 1178 static inline __device__ __host__ \ 1179 type make_##type(comp x, comp y, comp z, comp w) { \ 1180 type r; r.x=x; r.y=y; r.z=z; r.w=w; \ 1185 DECLOP_MAKE_ONE_COMPONENT(
unsigned char, uchar1);
1186 DECLOP_MAKE_TWO_COMPONENT(
unsigned char, uchar2);
1187 DECLOP_MAKE_THREE_COMPONENT(
unsigned char, uchar3);
1188 DECLOP_MAKE_FOUR_COMPONENT(
unsigned char, uchar4);
1190 DECLOP_MAKE_ONE_COMPONENT(
signed char, char1);
1191 DECLOP_MAKE_TWO_COMPONENT(
signed char, char2);
1192 DECLOP_MAKE_THREE_COMPONENT(
signed char, char3);
1193 DECLOP_MAKE_FOUR_COMPONENT(
signed char, char4);
1195 DECLOP_MAKE_ONE_COMPONENT(
unsigned short, ushort1);
1196 DECLOP_MAKE_TWO_COMPONENT(
unsigned short, ushort2);
1197 DECLOP_MAKE_THREE_COMPONENT(
unsigned short, ushort3);
1198 DECLOP_MAKE_FOUR_COMPONENT(
unsigned short, ushort4);
1200 DECLOP_MAKE_ONE_COMPONENT(
signed short, short1);
1201 DECLOP_MAKE_TWO_COMPONENT(
signed short, short2);
1202 DECLOP_MAKE_THREE_COMPONENT(
signed short, short3);
1203 DECLOP_MAKE_FOUR_COMPONENT(
signed short, short4);
1205 DECLOP_MAKE_ONE_COMPONENT(
unsigned int, uint1);
1206 DECLOP_MAKE_TWO_COMPONENT(
unsigned int, uint2);
1207 DECLOP_MAKE_THREE_COMPONENT(
unsigned int, uint3);
1208 DECLOP_MAKE_FOUR_COMPONENT(
unsigned int, uint4);
1210 DECLOP_MAKE_ONE_COMPONENT(
signed int, int1);
1211 DECLOP_MAKE_TWO_COMPONENT(
signed int, int2);
1212 DECLOP_MAKE_THREE_COMPONENT(
signed int, int3);
1213 DECLOP_MAKE_FOUR_COMPONENT(
signed int, int4);
1215 DECLOP_MAKE_ONE_COMPONENT(
float, float1);
1216 DECLOP_MAKE_TWO_COMPONENT(
float, float2);
1217 DECLOP_MAKE_THREE_COMPONENT(
float, float3);
1218 DECLOP_MAKE_FOUR_COMPONENT(
float, float4);
1220 DECLOP_MAKE_ONE_COMPONENT(
double, double1);
1221 DECLOP_MAKE_TWO_COMPONENT(
double, double2);
1222 DECLOP_MAKE_THREE_COMPONENT(
double, double3);
1223 DECLOP_MAKE_FOUR_COMPONENT(
double, double4);
1225 DECLOP_MAKE_ONE_COMPONENT(
unsigned long, ulong1);
1226 DECLOP_MAKE_TWO_COMPONENT(
unsigned long, ulong2);
1227 DECLOP_MAKE_THREE_COMPONENT(
unsigned long, ulong3);
1228 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long, ulong4);
1230 DECLOP_MAKE_ONE_COMPONENT(
signed long, long1);
1231 DECLOP_MAKE_TWO_COMPONENT(
signed long, long2);
1232 DECLOP_MAKE_THREE_COMPONENT(
signed long, long3);
1233 DECLOP_MAKE_FOUR_COMPONENT(
signed long, long4);
1235 DECLOP_MAKE_ONE_COMPONENT(
unsigned long long, ulonglong1);
1236 DECLOP_MAKE_TWO_COMPONENT(
unsigned long long, ulonglong2);
1237 DECLOP_MAKE_THREE_COMPONENT(
unsigned long long, ulonglong3);
1238 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long long, ulonglong4);
1240 DECLOP_MAKE_ONE_COMPONENT(
signed long long, longlong1);
1241 DECLOP_MAKE_TWO_COMPONENT(
signed long long, longlong2);
1242 DECLOP_MAKE_THREE_COMPONENT(
signed long long, longlong3);
1243 DECLOP_MAKE_FOUR_COMPONENT(
signed long long, longlong4);
1244 #else // defined(_MSC_VER) 1245 #include <mmintrin.h> 1246 #include <xmmintrin.h> 1247 #include <emmintrin.h> 1248 #include <immintrin.h> 1250 typedef union {
char data; } char1;
1251 typedef union {
char data[2]; } char2;
1252 typedef union {
char data[4]; } char4;
1253 typedef union { char4 data; } char3;
1254 typedef union { __m64 data; } char8;
1255 typedef union { __m128i data; } char16;
1257 typedef union {
unsigned char data; } uchar1;
1258 typedef union {
unsigned char data[2]; } uchar2;
1259 typedef union {
unsigned char data[4]; } uchar4;
1260 typedef union { uchar4 data; } uchar3;
1261 typedef union { __m64 data; } uchar8;
1262 typedef union { __m128i data; } uchar16;
1264 typedef union {
short data; } short1;
1265 typedef union {
short data[2]; } short2;
1266 typedef union { __m64 data; } short4;
1267 typedef union { short4 data; } short3;
1268 typedef union { __m128i data; } short8;
1269 typedef union { __m128i data[2]; } short16;
1271 typedef union {
unsigned short data; } ushort1;
1272 typedef union {
unsigned short data[2]; } ushort2;
1273 typedef union { __m64 data; } ushort4;
1274 typedef union { ushort4 data; } ushort3;
1275 typedef union { __m128i data; } ushort8;
1276 typedef union { __m128i data[2]; } ushort16;
1278 typedef union {
int data; } int1;
1279 typedef union { __m64 data; } int2;
1280 typedef union { __m128i data; } int4;
1281 typedef union { int4 data; } int3;
1282 typedef union { __m128i data[2]; } int8;
1283 typedef union { __m128i data[4];} int16;
1285 typedef union {
unsigned int data; } uint1;
1286 typedef union { __m64 data; } uint2;
1287 typedef union { __m128i data; } uint4;
1288 typedef union { uint4 data; } uint3;
1289 typedef union { __m128i data[2]; } uint8;
1290 typedef union { __m128i data[4]; } uint16;
1292 #if !defined(_WIN64) 1293 typedef union {
int data; } long1;
1294 typedef union { __m64 data; } long2;
1295 typedef union { __m128i data; } long4;
1296 typedef union { long4 data; } long3;
1297 typedef union { __m128i data[2]; } long8;
1298 typedef union { __m128i data[4]; } long16;
1300 typedef union {
unsigned int data; } ulong1;
1301 typedef union { __m64 data; } ulong2;
1302 typedef union { __m128i data; } ulong4;
1303 typedef union { ulong4 data; } ulong3;
1304 typedef union { __m128i data[2]; } ulong8;
1305 typedef union { __m128i data[4]; } ulong16;
1306 #else // defined(_WIN64) 1307 typedef union { __m64 data; } long1;
1308 typedef union { __m128i data; } long2;
1309 typedef union { __m128i data[2]; } long4;
1310 typedef union { long4 data; } long3;
1311 typedef union { __m128i data[4]; } long8;
1312 typedef union { __m128i data[8]; } long16;
1314 typedef union { __m64 data; } ulong1;
1315 typedef union { __m128i data; } ulong2;
1316 typedef union { __m128i data[2]; } ulong4;
1317 typedef union { ulong4 data; } ulong3;
1318 typedef union { __m128i data[4]; } ulong8;
1319 typedef union { __m128i data[8]; } ulong16;
1320 #endif // defined(_WIN64) 1322 typedef union { __m64 data; } longlong1;
1323 typedef union { __m128i data; } longlong2;
1324 typedef union { __m128i data[2]; } longlong4;
1325 typedef union { longlong4 data; } longlong3;
1326 typedef union { __m128i data[4]; } longlong8;
1327 typedef union { __m128i data[8]; } longlong16;
1329 typedef union { __m64 data; } ulonglong1;
1330 typedef union { __m128i data; } ulonglong2;
1331 typedef union { __m128i data[2]; } ulonglong4;
1332 typedef union { ulonglong4 data; } ulonglong3;
1333 typedef union { __m128i data[4]; } ulonglong8;
1334 typedef union { __m128i data[8]; } ulonglong16;
1336 typedef union {
float data; } float1;
1337 typedef union { __m64 data; } float2;
1338 typedef union { __m128 data; } float4;
1339 typedef union { float4 data; } float3;
1340 typedef union { __m256 data; } float8;
1341 typedef union { __m256 data[2]; } float16;
1343 typedef union {
double data; } double1;
1344 typedef union { __m128d data; } double2;
1345 typedef union { __m256d data; } double4;
1346 typedef union { double4 data; } double3;
1347 typedef union { __m256d data[2]; } double8;
1348 typedef union { __m256d data[4]; } double16;
1350 #endif // defined(_MSC_VER)
#define __host__
Definition: host_defines.h:41
Definition: hip_runtime.h:202