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) 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]; }
117 typename std::enable_if<
118 !std::is_same<U, T>{} &&
121 T,
typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* =
nullptr>
123 operator U() const noexcept {
return static_cast<U
>(data[idx]); }
126 typename std::enable_if<
127 !std::is_same<U, T>{} &&
130 T,
typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* =
nullptr>
132 operator U() const volatile noexcept {
return static_cast<U
>(data[idx]); }
135 operator T&() noexcept {
136 return reinterpret_cast< 137 T (&)[sizeof(Vector) / sizeof(T)]
>(data)[idx];
140 operator volatile T&()
volatile noexcept {
141 return reinterpret_cast< 142 volatile T (&)[sizeof(Vector) / sizeof(T)]
>(data)[idx];
146 Address operator&() const noexcept {
return Address{
this}; }
149 Scalar_accessor& operator=(
const Scalar_accessor& x) noexcept {
150 data[idx] = x.data[idx];
155 Scalar_accessor& operator=(T x) noexcept {
161 volatile Scalar_accessor& operator=(T x)
volatile noexcept {
168 Scalar_accessor& operator++() noexcept {
173 T operator++(
int) noexcept {
179 Scalar_accessor& operator--() noexcept {
184 T operator--(
int) noexcept {
194 typename std::enable_if<
195 std::is_convertible<U, T>{}>::type* =
nullptr>
197 Scalar_accessor& operator+=(U x) noexcept {
203 typename std::enable_if<
204 std::is_convertible<U, T>{}>::type* =
nullptr>
206 Scalar_accessor& operator-=(U x) noexcept {
213 typename std::enable_if<
214 std::is_convertible<U, T>{}>::type* =
nullptr>
216 Scalar_accessor& operator*=(U x) noexcept {
222 typename std::enable_if<
223 std::is_convertible<U, T>{}>::type* =
nullptr>
225 Scalar_accessor& operator/=(U x) noexcept {
231 typename std::enable_if<std::is_convertible<U, T>{} &&
232 std::is_integral<U>{}>::type* =
nullptr>
234 Scalar_accessor& operator%=(U x) noexcept {
241 typename std::enable_if<std::is_convertible<U, T>{} &&
242 std::is_integral<U>{}>::type* =
nullptr>
244 Scalar_accessor& operator>>=(U x) noexcept {
250 typename std::enable_if<std::is_convertible<U, T>{} &&
251 std::is_integral<U>{}>::type* =
nullptr>
253 Scalar_accessor& operator<<=(U x) noexcept {
259 typename std::enable_if<std::is_convertible<U, T>{} &&
260 std::is_integral<U>{}>::type* =
nullptr>
262 Scalar_accessor& operator&=(U x) noexcept {
268 typename std::enable_if<std::is_convertible<U, T>{} &&
269 std::is_integral<U>{}>::type* =
nullptr>
271 Scalar_accessor& operator|=(U x) noexcept {
277 typename std::enable_if<std::is_convertible<U, T>{} &&
278 std::is_integral<U>{}>::type* =
nullptr>
280 Scalar_accessor& operator^=(U x) noexcept {
287 template<
typename T,
unsigned int n>
struct HIP_vector_base;
290 struct HIP_vector_base<T, 1> {
291 using Native_vec_ = T __NATIVE_VECTOR__(1, T);
295 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
298 using value_type = T;
301 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
302 #if __has_attribute(ext_vector_type) 313 struct HIP_vector_base<T, 2> {
314 using Native_vec_ = T __NATIVE_VECTOR__(2, T);
318 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
319 hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
322 using value_type = T;
325 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
326 #if __has_attribute(ext_vector_type) 338 struct HIP_vector_base<T, 3> {
344 Native_vec_() =
default;
348 Native_vec_(T x) noexcept : d{x, x, x} {}
351 Native_vec_(T x, T y, T z) noexcept : d{x, y, z} {}
354 Native_vec_(
const Native_vec_&) =
default;
357 Native_vec_(Native_vec_&&) =
default;
359 ~Native_vec_() =
default;
362 Native_vec_& operator=(
const Native_vec_&) =
default;
364 Native_vec_& operator=(Native_vec_&&) =
default;
367 T& operator[](
unsigned int idx) noexcept {
return d[idx]; }
369 T operator[](
unsigned int idx)
const noexcept {
return d[idx]; }
372 Native_vec_& operator+=(
const Native_vec_& x) noexcept
374 for (
auto i = 0u; i != 3u; ++i) d[i] += x.d[i];
378 Native_vec_& operator-=(
const Native_vec_& x) noexcept
380 for (
auto i = 0u; i != 3u; ++i) d[i] -= x.d[i];
385 Native_vec_& operator*=(
const Native_vec_& x) noexcept
387 for (
auto i = 0u; i != 3u; ++i) d[i] *= x.d[i];
391 Native_vec_& operator/=(
const Native_vec_& x) noexcept
393 for (
auto i = 0u; i != 3u; ++i) d[i] /= x.d[i];
399 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
401 Native_vec_ operator-() const noexcept
404 for (
auto&& x : r.d) x = -x;
410 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
412 Native_vec_ operator~() const noexcept
415 for (
auto&& x : r.d) x = ~x;
420 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
422 Native_vec_& operator%=(
const Native_vec_& x) noexcept
424 for (
auto i = 0u; i != 3u; ++i) d[i] %= x.d[i];
429 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
431 Native_vec_& operator^=(
const Native_vec_& x) noexcept
433 for (
auto i = 0u; i != 3u; ++i) d[i] ^= x.d[i];
438 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
440 Native_vec_& operator|=(
const Native_vec_& x) noexcept
442 for (
auto i = 0u; i != 3u; ++i) d[i] |= x.d[i];
447 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
449 Native_vec_& operator&=(
const Native_vec_& x) noexcept
451 for (
auto i = 0u; i != 3u; ++i) d[i] &= x.d[i];
456 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
458 Native_vec_& operator>>=(
const Native_vec_& x) noexcept
460 for (
auto i = 0u; i != 3u; ++i) d[i] >>= x.d[i];
465 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
467 Native_vec_& operator<<=(
const Native_vec_& x) noexcept
469 for (
auto i = 0u; i != 3u; ++i) d[i] <<= x.d[i];
473 using Vec3_cmp =
int __attribute__((vector_size(4 *
sizeof(
int))));
475 Vec3_cmp operator==(
const Native_vec_& x)
const noexcept
477 return Vec3_cmp{d[0] == x.d[0], d[1] == x.d[1], d[2] == x.d[2]};
490 using value_type = T;
494 struct HIP_vector_base<T, 4> {
495 using Native_vec_ = T __NATIVE_VECTOR__(4, T);
499 hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
500 hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
501 hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
502 hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
505 using value_type = T;
508 HIP_vector_base& operator=(
const HIP_vector_base& x) noexcept {
509 #if __has_attribute(ext_vector_type) 522 template<
typename T,
unsigned int rank>
523 struct HIP_vector_type :
public HIP_vector_base<T, rank> {
524 using HIP_vector_base<T, rank>::data;
525 using typename HIP_vector_base<T, rank>::Native_vec_;
528 HIP_vector_type() =
default;
531 typename std::enable_if<
532 std::is_convertible<U, T>{}>::type* =
nullptr>
534 HIP_vector_type(U x) noexcept
536 for (
auto i = 0u; i != rank; ++i) data[i] = x;
540 typename std::enable_if<
541 (rank > 1) &&
sizeof...(Us) == rank>::type* =
nullptr>
543 HIP_vector_type(Us... xs) noexcept
545 #if __has_attribute(ext_vector_type) 546 new (&data) Native_vec_{
static_cast<T
>(xs)...};
548 new (&data) std::array<T, rank>{
static_cast<T
>(xs)...};
552 HIP_vector_type(
const HIP_vector_type&) =
default;
554 HIP_vector_type(HIP_vector_type&&) =
default;
556 ~HIP_vector_type() =
default;
559 HIP_vector_type& operator=(
const HIP_vector_type&) =
default;
561 HIP_vector_type& operator=(HIP_vector_type&&) =
default;
565 HIP_vector_type& operator++() noexcept
567 return *
this += HIP_vector_type{1};
570 HIP_vector_type operator++(
int) noexcept
578 HIP_vector_type& operator--() noexcept
580 return *
this -= HIP_vector_type{1};
583 HIP_vector_type operator--(
int) noexcept
591 HIP_vector_type& operator+=(
const HIP_vector_type& x) noexcept
598 typename std::enable_if<
599 std::is_convertible<U, T>{}>::type* =
nullptr>
601 HIP_vector_type& operator+=(U x) noexcept
603 return *
this += HIP_vector_type{x};
607 HIP_vector_type& operator-=(
const HIP_vector_type& x) noexcept
614 typename std::enable_if<
615 std::is_convertible<U, T>{}>::type* =
nullptr>
617 HIP_vector_type& operator-=(U x) noexcept
619 return *
this -= HIP_vector_type{x};
623 HIP_vector_type& operator*=(
const HIP_vector_type& x) noexcept
630 typename std::enable_if<
631 std::is_convertible<U, T>{}>::type* =
nullptr>
633 HIP_vector_type& operator*=(U x) noexcept
635 return *
this *= HIP_vector_type{x};
639 HIP_vector_type& operator/=(
const HIP_vector_type& x) noexcept
646 typename std::enable_if<
647 std::is_convertible<U, T>{}>::type* =
nullptr>
649 HIP_vector_type& operator/=(U x) noexcept
651 return *
this /= HIP_vector_type{x};
656 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
658 HIP_vector_type operator-() noexcept
661 tmp.data = -tmp.data;
667 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
669 HIP_vector_type operator~() noexcept
671 HIP_vector_type r{*
this};
678 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
680 HIP_vector_type& operator%=(
const HIP_vector_type& x) noexcept
688 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
690 HIP_vector_type& operator^=(
const HIP_vector_type& x) noexcept
698 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
700 HIP_vector_type& operator|=(
const HIP_vector_type& x) noexcept
708 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
710 HIP_vector_type& operator&=(
const HIP_vector_type& x) noexcept
718 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
720 HIP_vector_type& operator>>=(
const HIP_vector_type& x) noexcept
728 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
730 HIP_vector_type& operator<<=(
const HIP_vector_type& x) noexcept
737 template<
typename T,
unsigned int n>
739 HIP_vector_type<T, n> operator+(
740 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
742 return HIP_vector_type<T, n>{x} += y;
744 template<
typename T,
unsigned int n,
typename U>
746 HIP_vector_type<T, n> operator+(
747 const HIP_vector_type<T, n>& x, U y) noexcept
749 return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
751 template<
typename T,
unsigned int n,
typename U>
753 HIP_vector_type<T, n> operator+(
754 U x,
const HIP_vector_type<T, n>& y) noexcept
756 return HIP_vector_type<T, n>{x} += y;
759 template<
typename T,
unsigned int n>
761 HIP_vector_type<T, n> operator-(
762 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
764 return HIP_vector_type<T, n>{x} -= y;
766 template<
typename T,
unsigned int n,
typename U>
768 HIP_vector_type<T, n> operator-(
769 const HIP_vector_type<T, n>& x, U y) noexcept
771 return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
773 template<
typename T,
unsigned int n,
typename U>
775 HIP_vector_type<T, n> operator-(
776 U x,
const HIP_vector_type<T, n>& y) noexcept
778 return HIP_vector_type<T, n>{x} -= y;
781 template<
typename T,
unsigned int n>
783 HIP_vector_type<T, n> operator*(
784 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
786 return HIP_vector_type<T, n>{x} *= y;
788 template<
typename T,
unsigned int n,
typename U>
790 HIP_vector_type<T, n> operator*(
791 const HIP_vector_type<T, n>& x, U y) noexcept
793 return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
795 template<
typename T,
unsigned int n,
typename U>
797 HIP_vector_type<T, n> operator*(
798 U x,
const HIP_vector_type<T, n>& y) noexcept
800 return HIP_vector_type<T, n>{x} *= y;
803 template<
typename T,
unsigned int n>
805 HIP_vector_type<T, n> operator/(
806 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
808 return HIP_vector_type<T, n>{x} /= y;
810 template<
typename T,
unsigned int n,
typename U>
812 HIP_vector_type<T, n> operator/(
813 const HIP_vector_type<T, n>& x, U y) noexcept
815 return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
817 template<
typename T,
unsigned int n,
typename U>
819 HIP_vector_type<T, n> operator/(
820 U x,
const HIP_vector_type<T, n>& y) noexcept
822 return HIP_vector_type<T, n>{x} /= y;
825 template<
typename T,
unsigned int n>
828 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
830 auto tmp = x.data == y.data;
831 for (
auto i = 0u; i != n; ++i)
if (tmp[i] == 0)
return false;
834 template<
typename T,
unsigned int n,
typename U>
836 bool operator==(
const HIP_vector_type<T, n>& x, U y) noexcept
838 return x == HIP_vector_type<T, n>{y};
840 template<
typename T,
unsigned int n,
typename U>
842 bool operator==(U x,
const HIP_vector_type<T, n>& y) noexcept
844 return HIP_vector_type<T, n>{x} == y;
847 template<
typename T,
unsigned int n>
850 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
854 template<
typename T,
unsigned int n,
typename U>
856 bool operator!=(
const HIP_vector_type<T, n>& x, U y) noexcept
860 template<
typename T,
unsigned int n,
typename U>
862 bool operator!=(U x,
const HIP_vector_type<T, n>& y) noexcept
870 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
872 HIP_vector_type<T, n> operator%(
873 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
875 return HIP_vector_type<T, n>{x} %= y;
881 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
883 HIP_vector_type<T, n> operator%(
884 const HIP_vector_type<T, n>& x, U y) noexcept
886 return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
892 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
894 HIP_vector_type<T, n> operator%(
895 U x,
const HIP_vector_type<T, n>& y) noexcept
897 return HIP_vector_type<T, n>{x} %= y;
903 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
905 HIP_vector_type<T, n> operator^(
906 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
908 return HIP_vector_type<T, n>{x} ^= y;
914 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
916 HIP_vector_type<T, n> operator^(
917 const HIP_vector_type<T, n>& x, U y) noexcept
919 return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
925 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
927 HIP_vector_type<T, n> operator^(
928 U x,
const HIP_vector_type<T, n>& y) noexcept
930 return HIP_vector_type<T, n>{x} ^= y;
936 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
938 HIP_vector_type<T, n> operator|(
939 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
941 return HIP_vector_type<T, n>{x} |= y;
947 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
949 HIP_vector_type<T, n> operator|(
950 const HIP_vector_type<T, n>& x, U y) noexcept
952 return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
958 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
960 HIP_vector_type<T, n> operator|(
961 U x,
const HIP_vector_type<T, n>& y) noexcept
963 return HIP_vector_type<T, n>{x} |= y;
969 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
971 HIP_vector_type<T, n> operator&(
972 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
974 return HIP_vector_type<T, n>{x} &= y;
980 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
982 HIP_vector_type<T, n> operator&(
983 const HIP_vector_type<T, n>& x, U y) noexcept
985 return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
991 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
993 HIP_vector_type<T, n> operator&(
994 U x,
const HIP_vector_type<T, n>& y) noexcept
996 return HIP_vector_type<T, n>{x} &= y;
1002 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1004 HIP_vector_type<T, n> operator>>(
1005 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1007 return HIP_vector_type<T, n>{x} >>= y;
1013 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1015 HIP_vector_type<T, n> operator>>(
1016 const HIP_vector_type<T, n>& x, U y) noexcept
1018 return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
1024 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1026 HIP_vector_type<T, n> operator>>(
1027 U x,
const HIP_vector_type<T, n>& y) noexcept
1029 return HIP_vector_type<T, n>{x} >>= y;
1035 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1037 HIP_vector_type<T, n> operator<<(
1038 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
1040 return HIP_vector_type<T, n>{x} <<= y;
1046 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1048 HIP_vector_type<T, n> operator<<(
1049 const HIP_vector_type<T, n>& x, U y) noexcept
1051 return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
1057 typename std::enable_if<std::is_arithmetic<U>::value>::type,
1058 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1060 HIP_vector_type<T, n> operator<<(
1061 U x,
const HIP_vector_type<T, n>& y) noexcept
1063 return HIP_vector_type<T, n>{x} <<= y;
1066 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 1067 using CUDA_name##1 = HIP_vector_type<T, 1>;\ 1068 using CUDA_name##2 = HIP_vector_type<T, 2>;\ 1069 using CUDA_name##3 = HIP_vector_type<T, 3>;\ 1070 using CUDA_name##4 = HIP_vector_type<T, 4>; 1072 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 1093 __MAKE_VECTOR_TYPE__(uchar,
unsigned char);
1094 __MAKE_VECTOR_TYPE__(
char,
char);
1095 __MAKE_VECTOR_TYPE__(ushort,
unsigned short);
1096 __MAKE_VECTOR_TYPE__(
short,
short);
1097 __MAKE_VECTOR_TYPE__(uint,
unsigned int);
1098 __MAKE_VECTOR_TYPE__(
int,
int);
1099 __MAKE_VECTOR_TYPE__(ulong,
unsigned long);
1100 __MAKE_VECTOR_TYPE__(
long,
long);
1101 __MAKE_VECTOR_TYPE__(ulonglong,
unsigned long long);
1102 __MAKE_VECTOR_TYPE__(longlong,
long long);
1103 __MAKE_VECTOR_TYPE__(
float,
float);
1104 __MAKE_VECTOR_TYPE__(
double,
double);
1107 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 1108 static inline __device__ __host__ \ 1109 type make_##type(comp x) { type r{x}; return r; } 1111 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 1112 static inline __device__ __host__ \ 1113 type make_##type(comp x, comp y) { type r{x, y}; return r; } 1115 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 1116 static inline __device__ __host__ \ 1117 type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; } 1119 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 1120 static inline __device__ __host__ \ 1121 type make_##type(comp x, comp y, comp z, comp w) { \ 1122 type r{x, y, z, w}; \ 1126 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 1127 static inline __device__ __host__ \ 1128 type make_##type(comp x) { type r; r.x =x; return r; } 1130 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 1131 static inline __device__ __host__ \ 1132 type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; } 1134 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 1135 static inline __device__ __host__ \ 1136 type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; } 1138 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 1139 static inline __device__ __host__ \ 1140 type make_##type(comp x, comp y, comp z, comp w) { \ 1141 type r; r.x=x; r.y=y; r.z=z; r.w=w; \ 1146 DECLOP_MAKE_ONE_COMPONENT(
unsigned char, uchar1);
1147 DECLOP_MAKE_TWO_COMPONENT(
unsigned char, uchar2);
1148 DECLOP_MAKE_THREE_COMPONENT(
unsigned char, uchar3);
1149 DECLOP_MAKE_FOUR_COMPONENT(
unsigned char, uchar4);
1151 DECLOP_MAKE_ONE_COMPONENT(
signed char, char1);
1152 DECLOP_MAKE_TWO_COMPONENT(
signed char, char2);
1153 DECLOP_MAKE_THREE_COMPONENT(
signed char, char3);
1154 DECLOP_MAKE_FOUR_COMPONENT(
signed char, char4);
1156 DECLOP_MAKE_ONE_COMPONENT(
unsigned short, ushort1);
1157 DECLOP_MAKE_TWO_COMPONENT(
unsigned short, ushort2);
1158 DECLOP_MAKE_THREE_COMPONENT(
unsigned short, ushort3);
1159 DECLOP_MAKE_FOUR_COMPONENT(
unsigned short, ushort4);
1161 DECLOP_MAKE_ONE_COMPONENT(
signed short, short1);
1162 DECLOP_MAKE_TWO_COMPONENT(
signed short, short2);
1163 DECLOP_MAKE_THREE_COMPONENT(
signed short, short3);
1164 DECLOP_MAKE_FOUR_COMPONENT(
signed short, short4);
1166 DECLOP_MAKE_ONE_COMPONENT(
unsigned int, uint1);
1167 DECLOP_MAKE_TWO_COMPONENT(
unsigned int, uint2);
1168 DECLOP_MAKE_THREE_COMPONENT(
unsigned int, uint3);
1169 DECLOP_MAKE_FOUR_COMPONENT(
unsigned int, uint4);
1171 DECLOP_MAKE_ONE_COMPONENT(
signed int, int1);
1172 DECLOP_MAKE_TWO_COMPONENT(
signed int, int2);
1173 DECLOP_MAKE_THREE_COMPONENT(
signed int, int3);
1174 DECLOP_MAKE_FOUR_COMPONENT(
signed int, int4);
1176 DECLOP_MAKE_ONE_COMPONENT(
float, float1);
1177 DECLOP_MAKE_TWO_COMPONENT(
float, float2);
1178 DECLOP_MAKE_THREE_COMPONENT(
float, float3);
1179 DECLOP_MAKE_FOUR_COMPONENT(
float, float4);
1181 DECLOP_MAKE_ONE_COMPONENT(
double, double1);
1182 DECLOP_MAKE_TWO_COMPONENT(
double, double2);
1183 DECLOP_MAKE_THREE_COMPONENT(
double, double3);
1184 DECLOP_MAKE_FOUR_COMPONENT(
double, double4);
1186 DECLOP_MAKE_ONE_COMPONENT(
unsigned long, ulong1);
1187 DECLOP_MAKE_TWO_COMPONENT(
unsigned long, ulong2);
1188 DECLOP_MAKE_THREE_COMPONENT(
unsigned long, ulong3);
1189 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long, ulong4);
1191 DECLOP_MAKE_ONE_COMPONENT(
signed long, long1);
1192 DECLOP_MAKE_TWO_COMPONENT(
signed long, long2);
1193 DECLOP_MAKE_THREE_COMPONENT(
signed long, long3);
1194 DECLOP_MAKE_FOUR_COMPONENT(
signed long, long4);
1196 DECLOP_MAKE_ONE_COMPONENT(
unsigned long long, ulonglong1);
1197 DECLOP_MAKE_TWO_COMPONENT(
unsigned long long, ulonglong2);
1198 DECLOP_MAKE_THREE_COMPONENT(
unsigned long long, ulonglong3);
1199 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long long, ulonglong4);
1201 DECLOP_MAKE_ONE_COMPONENT(
signed long long, longlong1);
1202 DECLOP_MAKE_TWO_COMPONENT(
signed long long, longlong2);
1203 DECLOP_MAKE_THREE_COMPONENT(
signed long long, longlong3);
1204 DECLOP_MAKE_FOUR_COMPONENT(
signed long long, longlong4);
1205 #else // defined(_MSC_VER) 1206 #include <mmintrin.h> 1207 #include <xmmintrin.h> 1208 #include <emmintrin.h> 1209 #include <immintrin.h> 1211 typedef union {
char data; } char1;
1212 typedef union {
char data[2]; } char2;
1213 typedef union {
char data[4]; } char4;
1214 typedef union { char4 data; } char3;
1215 typedef union { __m64 data; } char8;
1216 typedef union { __m128i data; } char16;
1218 typedef union {
unsigned char data; } uchar1;
1219 typedef union {
unsigned char data[2]; } uchar2;
1220 typedef union {
unsigned char data[4]; } uchar4;
1221 typedef union { uchar4 data; } uchar3;
1222 typedef union { __m64 data; } uchar8;
1223 typedef union { __m128i data; } uchar16;
1225 typedef union {
short data; } short1;
1226 typedef union {
short data[2]; } short2;
1227 typedef union { __m64 data; } short4;
1228 typedef union { short4 data; } short3;
1229 typedef union { __m128i data; } short8;
1230 typedef union { __m128i data[2]; } short16;
1232 typedef union {
unsigned short data; } ushort1;
1233 typedef union {
unsigned short data[2]; } ushort2;
1234 typedef union { __m64 data; } ushort4;
1235 typedef union { ushort4 data; } ushort3;
1236 typedef union { __m128i data; } ushort8;
1237 typedef union { __m128i data[2]; } ushort16;
1239 typedef union {
int data; } int1;
1240 typedef union { __m64 data; } int2;
1241 typedef union { __m128i data; } int4;
1242 typedef union { int4 data; } int3;
1243 typedef union { __m128i data[2]; } int8;
1244 typedef union { __m128i data[4];} int16;
1246 typedef union {
unsigned int data; } uint1;
1247 typedef union { __m64 data; } uint2;
1248 typedef union { __m128i data; } uint4;
1249 typedef union { uint4 data; } uint3;
1250 typedef union { __m128i data[2]; } uint8;
1251 typedef union { __m128i data[4]; } uint16;
1253 #if !defined(_WIN64) 1254 typedef union {
int data; } long1;
1255 typedef union { __m64 data; } long2;
1256 typedef union { __m128i data; } long4;
1257 typedef union { long4 data; } long3;
1258 typedef union { __m128i data[2]; } long8;
1259 typedef union { __m128i data[4]; } long16;
1261 typedef union {
unsigned int data; } ulong1;
1262 typedef union { __m64 data; } ulong2;
1263 typedef union { __m128i data; } ulong4;
1264 typedef union { ulong4 data; } ulong3;
1265 typedef union { __m128i data[2]; } ulong8;
1266 typedef union { __m128i data[4]; } ulong16;
1267 #else // defined(_WIN64) 1268 typedef union { __m64 data; } long1;
1269 typedef union { __m128i data; } long2;
1270 typedef union { __m128i data[2]; } long4;
1271 typedef union { long4 data; } long3;
1272 typedef union { __m128i data[4]; } long8;
1273 typedef union { __m128i data[8]; } long16;
1275 typedef union { __m64 data; } ulong1;
1276 typedef union { __m128i data; } ulong2;
1277 typedef union { __m128i data[2]; } ulong4;
1278 typedef union { ulong4 data; } ulong3;
1279 typedef union { __m128i data[4]; } ulong8;
1280 typedef union { __m128i data[8]; } ulong16;
1281 #endif // defined(_WIN64) 1283 typedef union { __m64 data; } longlong1;
1284 typedef union { __m128i data; } longlong2;
1285 typedef union { __m128i data[2]; } longlong4;
1286 typedef union { longlong4 data; } longlong3;
1287 typedef union { __m128i data[4]; } longlong8;
1288 typedef union { __m128i data[8]; } longlong16;
1290 typedef union { __m64 data; } ulonglong1;
1291 typedef union { __m128i data; } ulonglong2;
1292 typedef union { __m128i data[2]; } ulonglong4;
1293 typedef union { ulonglong4 data; } ulonglong3;
1294 typedef union { __m128i data[4]; } ulonglong8;
1295 typedef union { __m128i data[8]; } ulonglong16;
1297 typedef union {
float data; } float1;
1298 typedef union { __m64 data; } float2;
1299 typedef union { __m128 data; } float4;
1300 typedef union { float4 data; } float3;
1301 typedef union { __m256 data; } float8;
1302 typedef union { __m256 data[2]; } float16;
1304 typedef union {
double data; } double1;
1305 typedef union { __m128d data; } double2;
1306 typedef union { __m256d data; } double4;
1307 typedef union { double4 data; } double3;
1308 typedef union { __m256d data[2]; } double8;
1309 typedef union { __m256d data[4]; } double16;
1311 #endif // defined(_MSC_VER)
#define __host__
Definition: host_defines.h:41
Definition: hip_runtime.h:202