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(__has_attribute) 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-() const noexcept
700 tmp.data = -tmp.data;
706 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
708 HIP_vector_type operator~() const 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(__has_attribute) 1246 #if defined(_MSC_VER) 1247 #include <mmintrin.h> 1248 #include <xmmintrin.h> 1249 #include <emmintrin.h> 1250 #include <immintrin.h> 1252 typedef union {
char data; }
char1;
1253 typedef union {
char data[2]; }
char2;
1254 typedef union {
char data[4]; }
char4;
1256 typedef union { __m64 data; }
char8;
1257 typedef union { __m128i data; }
char16;
1259 typedef union {
unsigned char data; }
uchar1;
1260 typedef union {
unsigned char data[2]; }
uchar2;
1261 typedef union {
unsigned char data[4]; }
uchar4;
1263 typedef union { __m64 data; }
uchar8;
1264 typedef union { __m128i data; }
uchar16;
1266 typedef union {
short data; }
short1;
1267 typedef union {
short data[2]; }
short2;
1268 typedef union { __m64 data; }
short4;
1270 typedef union { __m128i data; }
short8;
1271 typedef union { __m128i data[2]; }
short16;
1273 typedef union {
unsigned short data; }
ushort1;
1274 typedef union {
unsigned short data[2]; }
ushort2;
1275 typedef union { __m64 data; }
ushort4;
1277 typedef union { __m128i data; }
ushort8;
1278 typedef union { __m128i data[2]; }
ushort16;
1280 typedef union {
int data; }
int1;
1281 typedef union { __m64 data; }
int2;
1282 typedef union { __m128i data; }
int4;
1284 typedef union { __m128i data[2]; }
int8;
1285 typedef union { __m128i data[4];}
int16;
1287 typedef union {
unsigned int data; }
uint1;
1288 typedef union { __m64 data; }
uint2;
1289 typedef union { __m128i data; }
uint4;
1291 typedef union { __m128i data[2]; }
uint8;
1292 typedef union { __m128i data[4]; }
uint16;
1294 #if !defined(_WIN64) 1295 typedef union {
int data; }
long1;
1296 typedef union { __m64 data; }
long2;
1297 typedef union { __m128i data; }
long4;
1299 typedef union { __m128i data[2]; }
long8;
1300 typedef union { __m128i data[4]; }
long16;
1302 typedef union {
unsigned int data; }
ulong1;
1303 typedef union { __m64 data; }
ulong2;
1304 typedef union { __m128i data; }
ulong4;
1306 typedef union { __m128i data[2]; }
ulong8;
1307 typedef union { __m128i data[4]; }
ulong16;
1308 #else // defined(_WIN64) 1309 typedef union { __m64 data; }
long1;
1310 typedef union { __m128i data; }
long2;
1311 typedef union { __m128i data[2]; }
long4;
1313 typedef union { __m128i data[4]; }
long8;
1314 typedef union { __m128i data[8]; }
long16;
1316 typedef union { __m64 data; }
ulong1;
1317 typedef union { __m128i data; }
ulong2;
1318 typedef union { __m128i data[2]; }
ulong4;
1320 typedef union { __m128i data[4]; }
ulong8;
1321 typedef union { __m128i data[8]; }
ulong16;
1322 #endif // defined(_WIN64) 1324 typedef union { __m64 data; }
longlong1;
1325 typedef union { __m128i data; }
longlong2;
1326 typedef union { __m128i data[2]; }
longlong4;
1328 typedef union { __m128i data[4]; }
longlong8;
1329 typedef union { __m128i data[8]; }
longlong16;
1333 typedef union { __m128i data[2]; }
ulonglong4;
1335 typedef union { __m128i data[4]; }
ulonglong8;
1338 typedef union {
float data; }
float1;
1339 typedef union { __m64 data; }
float2;
1340 typedef union { __m128 data; }
float4;
1342 typedef union { __m256 data; }
float8;
1343 typedef union { __m256 data[2]; }
float16;
1345 typedef union {
double data; }
double1;
1346 typedef union { __m128d data; }
double2;
1347 typedef union { __m256d data; }
double4;
1349 typedef union { __m256d data[2]; }
double8;
1350 typedef union { __m256d data[4]; }
double16;
1352 #else // !defined(_MSC_VER) 1361 typedef union {
unsigned char data; }
uchar1;
1362 typedef union {
unsigned char data[2]; }
uchar2;
1363 typedef union {
unsigned char data[4]; }
uchar4;
1364 typedef union {
unsigned char data[8]; }
uchar8;
1365 typedef union {
unsigned char data[16]; }
uchar16;
1376 typedef union {
unsigned short data[2]; }
ushort2;
1377 typedef union {
unsigned short data[4]; }
ushort4;
1378 typedef union {
unsigned short data[8]; }
ushort8;
1383 typedef union {
int data[2]; }
int2;
1384 typedef union {
int data[4]; }
int4;
1385 typedef union {
int data[8]; }
int8;
1389 typedef union {
unsigned int data; }
uint1;
1390 typedef union {
unsigned int data[2]; }
uint2;
1391 typedef union {
unsigned int data[4]; }
uint4;
1392 typedef union {
unsigned int data[8]; }
uint8;
1393 typedef union {
unsigned int data[16]; }
uint16;
1403 typedef union {
unsigned long data; }
ulong1;
1404 typedef union {
unsigned long data[2]; }
ulong2;
1405 typedef union {
unsigned long data[4]; }
ulong4;
1406 typedef union {
unsigned long data[8]; }
ulong8;
1407 typedef union {
unsigned long data[16]; }
ulong16;
1438 #endif // defined(_MSC_VER) 1439 #endif // defined(__has_attribute) Definition: hip_vector_types.h:1391
Definition: hip_vector_types.h:1390
Definition: hip_vector_types.h:1377
Definition: hip_vector_types.h:1418
Definition: hip_vector_types.h:1362
Definition: hip_vector_types.h:1369
Definition: hip_vector_types.h:1424
Definition: hip_vector_types.h:1422
Definition: hip_vector_types.h:1378
Definition: hip_vector_types.h:1432
Definition: hip_vector_types.h:1384
Definition: hip_vector_types.h:1394
Definition: hip_vector_types.h:1358
Definition: hip_vector_types.h:1389
Definition: hip_vector_types.h:1408
Definition: hip_vector_types.h:1386
Definition: hip_vector_types.h:1415
Definition: hip_vector_types.h:1417
Definition: hip_vector_types.h:1385
Definition: hip_vector_types.h:1406
Definition: hip_vector_types.h:1357
Definition: hip_vector_types.h:1431
Definition: hip_vector_types.h:1414
Definition: hip_vector_types.h:1387
Definition: hip_vector_types.h:1403
#define __host__
Definition: host_defines.h:41
Definition: hip_vector_types.h:1434
Definition: hip_vector_types.h:1405
Definition: hip_vector_types.h:1356
Definition: hip_vector_types.h:1412
Definition: hip_vector_types.h:1355
Definition: hip_vector_types.h:1365
Definition: hip_vector_types.h:1372
Definition: hip_vector_types.h:1393
Definition: hip_vector_types.h:1399
Definition: hip_vector_types.h:1433
Definition: hip_vector_types.h:1427
Definition: hip_vector_types.h:1413
Definition: hip_vector_types.h:1383
Definition: hip_vector_types.h:1404
Definition: hip_vector_types.h:1375
Definition: hip_vector_types.h:1407
Definition: hip_vector_types.h:1359
Definition: hip_vector_types.h:1410
Definition: hip_vector_types.h:1366
Definition: hip_vector_types.h:1354
Definition: hip_vector_types.h:1398
Definition: hip_vector_types.h:1373
Definition: hip_vector_types.h:1397
Definition: hip_vector_types.h:1379
Definition: hip_vector_types.h:1426
Definition: hip_runtime.h:210
Definition: hip_vector_types.h:1420
Definition: hip_vector_types.h:1364
Definition: hip_vector_types.h:1380
Definition: hip_vector_types.h:1436
Definition: hip_vector_types.h:1371
Definition: hip_vector_types.h:1382
Definition: hip_vector_types.h:1392
Definition: hip_vector_types.h:1376
Definition: hip_vector_types.h:1400
Definition: hip_vector_types.h:1361
Definition: hip_vector_types.h:1368
Definition: hip_vector_types.h:1425
Definition: hip_vector_types.h:1419
Definition: hip_vector_types.h:1411
Definition: hip_vector_types.h:1421
Definition: hip_vector_types.h:1363
Definition: hip_vector_types.h:1435
Definition: hip_vector_types.h:1401
Definition: hip_vector_types.h:1370
Definition: hip_vector_types.h:1396
Definition: hip_vector_types.h:1429
Definition: hip_vector_types.h:1428