28#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_VECTOR_TYPES_H
29#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_VECTOR_TYPES_H
33#if defined(__HIPCC_RTC__)
34 #define __HOST_DEVICE__ __device__
36 #define __HOST_DEVICE__ __host__ __device__
39#if defined(__has_attribute)
40 #if __has_attribute(ext_vector_type)
41 #define __NATIVE_VECTOR__(n, T) T __attribute__((ext_vector_type(n)))
43 #define __NATIVE_VECTOR__(n, T) T[n]
46#if defined(__cplusplus)
47#if !defined(__HIPCC_RTC__)
50 #include <type_traits>
55template <
class _Tp, _Tp __v>
struct integral_constant {
56 static constexpr const _Tp value = __v;
57 typedef _Tp value_type;
58 typedef integral_constant type;
59 constexpr operator value_type()
const {
return value; }
60 constexpr value_type operator()()
const {
return value; }
62template <
class _Tp, _Tp __v>
constexpr const _Tp integral_constant<_Tp, __v>::value;
64typedef integral_constant<bool, true> true_type;
65typedef integral_constant<bool, false> false_type;
67template <
bool B>
using bool_constant = integral_constant<bool, B>;
68typedef bool_constant<true> true_type;
69typedef bool_constant<false> false_type;
71template <
bool __B,
class __T =
void>
struct enable_if {};
72template <
class __T>
struct enable_if<true, __T> {
typedef __T type; };
74template<
bool _B>
struct true_or_false_type :
public false_type {};
75template<>
struct true_or_false_type<true> :
public true_type {};
77template <
class _Tp>
struct is_integral :
public false_type {};
78template <>
struct is_integral<bool> :
public true_type {};
79template <>
struct is_integral<char> :
public true_type {};
80template <>
struct is_integral<signed char> :
public true_type {};
81template <>
struct is_integral<unsigned char> :
public true_type {};
82template <>
struct is_integral<wchar_t> :
public true_type {};
83template <>
struct is_integral<short> :
public true_type {};
84template <>
struct is_integral<unsigned short> :
public true_type {};
85template <>
struct is_integral<int> :
public true_type {};
86template <>
struct is_integral<unsigned int> :
public true_type {};
87template <>
struct is_integral<long> :
public true_type {};
88template <>
struct is_integral<unsigned long> :
public true_type {};
89template <>
struct is_integral<long long> :
public true_type {};
90template <>
struct is_integral<unsigned long long> :
public true_type {};
92template <
class _Tp>
struct is_arithmetic :
public false_type {};
93template <>
struct is_arithmetic<bool> :
public true_type {};
94template <>
struct is_arithmetic<char> :
public true_type {};
95template <>
struct is_arithmetic<signed char> :
public true_type {};
96template <>
struct is_arithmetic<unsigned char> :
public true_type {};
97template <>
struct is_arithmetic<wchar_t> :
public true_type {};
98template <>
struct is_arithmetic<short> :
public true_type {};
99template <>
struct is_arithmetic<unsigned short> :
public true_type {};
100template <>
struct is_arithmetic<int> :
public true_type {};
101template <>
struct is_arithmetic<unsigned int> :
public true_type {};
102template <>
struct is_arithmetic<long> :
public true_type {};
103template <>
struct is_arithmetic<unsigned long> :
public true_type {};
104template <>
struct is_arithmetic<long long> :
public true_type {};
105template <>
struct is_arithmetic<unsigned long long> :
public true_type {};
106template <>
struct is_arithmetic<float> :
public true_type {};
107template <>
struct is_arithmetic<double> :
public true_type {};
109template<
typename _Tp>
struct is_floating_point :
public false_type {};
110template<>
struct is_floating_point<float> :
public true_type {};
111template<>
struct is_floating_point<double> :
public true_type {};
112template<>
struct is_floating_point<long double> :
public true_type {};
114template <
typename __T,
typename __U>
struct is_same :
public false_type {};
115template <
typename __T>
struct is_same<__T, __T> :
public true_type {};
117template<typename _Tp, bool = is_arithmetic<_Tp>::value>
118 struct is_signed :
public false_type {};
119template<
typename _Tp>
120 struct is_signed<_Tp, true> :
public true_or_false_type<_Tp(-1) < _Tp(0)> {};
122template <class _T1, class _T2> struct is_convertible
123 : public true_or_false_type<__is_convertible_to(_T1, _T2)> {};
125template<typename _CharT> struct char_traits;
126template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_istream;
127template<typename _CharT, typename _Traits = char_traits<_CharT>> class basic_ostream;
128typedef basic_istream<char> istream;
129typedef basic_ostream<char> ostream;
131template <typename __T> struct is_scalar : public integral_constant<bool, __is_scalar(__T)> {};
138 unsigned int next_pot(unsigned int x) {
140 return 1u << (32u - __builtin_clz(x - 1u));
144 template<typename T, unsigned int n> struct HIP_vector_base;
147 struct HIP_vector_base<T, 1> {
148 using Native_vec_ = __NATIVE_VECTOR__(1, T);
157 using value_type = T;
160 HIP_vector_base() = default;
164 HIP_vector_base(T x_) noexcept : data{x_} {}
167 HIP_vector_base(const HIP_vector_base&) = default;
170 HIP_vector_base(HIP_vector_base&&) = default;
172 ~HIP_vector_base() = default;
174 HIP_vector_base& operator=(const HIP_vector_base&) = default;
178 struct HIP_vector_base<T, 2> {
179 using Native_vec_ = __NATIVE_VECTOR__(2, T);
182 #if !__has_attribute(ext_vector_type)
183 alignas(hip_impl::next_pot(2 * sizeof(T)))
193 using value_type = T;
196 HIP_vector_base() = default;
200 HIP_vector_base(T x_) noexcept : data{x_, x_} {}
203 HIP_vector_base(T x_, T y_) noexcept : data{x_, y_} {}
206 HIP_vector_base(const HIP_vector_base&) = default;
209 HIP_vector_base(HIP_vector_base&&) = default;
211 ~HIP_vector_base() = default;
213 HIP_vector_base& operator=(const HIP_vector_base&) = default;
217 struct HIP_vector_base<T, 3> {
222 Native_vec_() = default;
227 Native_vec_(T x_) noexcept : d{x_, x_, x_} {}
230 Native_vec_(T x_, T y_, T z_) noexcept : d{x_, y_, z_} {}
233 Native_vec_(const Native_vec_&) = default;
236 Native_vec_(Native_vec_&&) = default;
238 ~Native_vec_() = default;
241 Native_vec_& operator=(const Native_vec_&) = default;
243 Native_vec_& operator=(Native_vec_&&) = default;
246 T& operator[](unsigned int idx) noexcept { return d[idx]; }
248 T operator[](unsigned int idx) const noexcept { return d[idx]; }
251 Native_vec_& operator+=(const Native_vec_& x_) noexcept
253 for (auto i = 0u; i != 3u; ++i) d[i] += x_.d[i];
257 Native_vec_& operator-=(const Native_vec_& x_) noexcept
259 for (auto i = 0u; i != 3u; ++i) d[i] -= x_.d[i];
264 Native_vec_& operator*=(const Native_vec_& x_) noexcept
266 for (auto i = 0u; i != 3u; ++i) d[i] *= x_.d[i];
270 Native_vec_& operator/=(const Native_vec_& x_) noexcept
272 for (auto i = 0u; i != 3u; ++i) d[i] /= x_.d[i];
278 typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
280 Native_vec_ operator-() const noexcept
283 for (auto&& x : r.d) x = -x;
289 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
291 Native_vec_ operator~() const noexcept
294 for (auto&& x : r.d) x = ~x;
299 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
301 Native_vec_& operator%=(const Native_vec_& x_) noexcept
303 for (auto i = 0u; i != 3u; ++i) d[i] %= x_.d[i];
308 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
310 Native_vec_& operator^=(const Native_vec_& x_) noexcept
312 for (auto i = 0u; i != 3u; ++i) d[i] ^= x_.d[i];
317 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
319 Native_vec_& operator|=(const Native_vec_& x_) noexcept
321 for (auto i = 0u; i != 3u; ++i) d[i] |= x_.d[i];
326 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
328 Native_vec_& operator&=(const Native_vec_& x_) noexcept
330 for (auto i = 0u; i != 3u; ++i) d[i] &= x_.d[i];
335 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
337 Native_vec_& operator>>=(const Native_vec_& x_) noexcept
339 for (auto i = 0u; i != 3u; ++i) d[i] >>= x_.d[i];
344 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
346 Native_vec_& operator<<=(const Native_vec_& x_) noexcept
348 for (auto i = 0u; i != 3u; ++i) d[i] <<= x_.d[i];
351#if defined (__INTEL_COMPILER)
355 using Vec3_cmp = _Vec3_cmp;
357 using Vec3_cmp = int __attribute__((vector_size(4 * sizeof(int))));
360 Vec3_cmp operator==(const Native_vec_& x_) const noexcept
362 return Vec3_cmp{d[0] == x_.d[0], d[1] == x_.d[1], d[2] == x_.d[2]};
375 using value_type = T;
378 HIP_vector_base() = default;
382 HIP_vector_base(T x_) noexcept : data{x_, x_, x_} {}
385 HIP_vector_base(T x_, T y_, T z_) noexcept : data{x_, y_, z_} {}
388 HIP_vector_base(const HIP_vector_base&) = default;
391 HIP_vector_base(HIP_vector_base&&) = default;
393 ~HIP_vector_base() = default;
396 HIP_vector_base& operator=(const HIP_vector_base&) = default;
398 HIP_vector_base& operator=(HIP_vector_base&&) = default;
402 struct HIP_vector_base<T, 4> {
403 using Native_vec_ = __NATIVE_VECTOR__(4, T);
406 #if !__has_attribute(ext_vector_type)
407 alignas(hip_impl::next_pot(4 * sizeof(T)))
419 using value_type = T;
422 HIP_vector_base() = default;
426 HIP_vector_base(T x_) noexcept : data{x_, x_, x_, x_} {}
429 HIP_vector_base(T x_, T y_, T z_, T w_) noexcept : data{x_, y_, z_, w_} {}
432 HIP_vector_base(const HIP_vector_base&) = default;
435 HIP_vector_base(HIP_vector_base&&) = default;
437 ~HIP_vector_base() = default;
439 HIP_vector_base& operator=(const HIP_vector_base&) = default;
442 template<typename T, unsigned int rank>
443 struct HIP_vector_type : public HIP_vector_base<T, rank> {
444 using HIP_vector_base<T, rank>::data;
445 using typename HIP_vector_base<T, rank>::Native_vec_;
448 HIP_vector_type() = default;
451 typename std::enable_if<
452 std::is_convertible<U, T>::value>::type* = nullptr>
456 HIP_vector_type(U x_) noexcept
457 : HIP_vector_base<T, rank>{static_cast<T>(x_)}
461 typename std::enable_if<
462 (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
465 HIP_vector_type(Us... xs) noexcept
466 : HIP_vector_base<T, rank>{static_cast<T>(xs)...}
470 HIP_vector_type(const HIP_vector_type&) = default;
473 HIP_vector_type(HIP_vector_type&&) = default;
475 ~HIP_vector_type() = default;
478 HIP_vector_type& operator=(const HIP_vector_type&) = default;
480 HIP_vector_type& operator=(HIP_vector_type&&) = default;
484 HIP_vector_type& operator++() noexcept
486 return *this += HIP_vector_type{1};
489 HIP_vector_type operator++(int) noexcept
497 HIP_vector_type& operator--() noexcept
499 return *this -= HIP_vector_type{1};
502 HIP_vector_type operator--(int) noexcept
510 HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
517 typename std::enable_if<
518 std::is_convertible<U, T>{}>::type* = nullptr>
520 HIP_vector_type& operator+=(U x) noexcept
522 return *this += HIP_vector_type{x};
526 HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
533 typename std::enable_if<
534 std::is_convertible<U, T>{}>::type* = nullptr>
536 HIP_vector_type& operator-=(U x) noexcept
538 return *this -= HIP_vector_type{x};
542 HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
548 friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
549 HIP_vector_type x, const HIP_vector_type& y) noexcept
551 return HIP_vector_type{ x } *= y;
556 typename std::enable_if<
557 std::is_convertible<U, T>{}>::type* = nullptr>
559 HIP_vector_type& operator*=(U x) noexcept
561 return *this *= HIP_vector_type{x};
564 friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator/(
565 HIP_vector_type x, const HIP_vector_type& y) noexcept
567 return HIP_vector_type{ x } /= y;
571 HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
578 typename std::enable_if<
579 std::is_convertible<U, T>{}>::type* = nullptr>
581 HIP_vector_type& operator/=(U x) noexcept
583 return *this /= HIP_vector_type{x};
588 typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
590 HIP_vector_type operator-() const noexcept
593 tmp.data = -tmp.data;
599 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
601 HIP_vector_type operator~() const noexcept
603 HIP_vector_type r{*this};
610 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
612 HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
620 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
622 HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
630 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
632 HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
640 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
642 HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
650 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
652 HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
660 typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
662 HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
669 template<typename T, unsigned int n>
673 HIP_vector_type<T, n> operator+(
674 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
676 return HIP_vector_type<T, n>{x} += y;
678 template<typename T, unsigned int n, typename U>
682 HIP_vector_type<T, n> operator+(
683 const HIP_vector_type<T, n>& x, U y) noexcept
685 return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
687 template<typename T, unsigned int n, typename U>
691 HIP_vector_type<T, n> operator+(
692 U x, const HIP_vector_type<T, n>& y) noexcept
694 return HIP_vector_type<T, n>{x} += y;
697 template<typename T, unsigned int n>
701 HIP_vector_type<T, n> operator-(
702 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
704 return HIP_vector_type<T, n>{x} -= y;
706 template<typename T, unsigned int n, typename U>
710 HIP_vector_type<T, n> operator-(
711 const HIP_vector_type<T, n>& x, U y) noexcept
713 return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
715 template<typename T, unsigned int n, typename U>
719 HIP_vector_type<T, n> operator-(
720 U x, const HIP_vector_type<T, n>& y) noexcept
722 return HIP_vector_type<T, n>{x} -= y;
725 template<typename T, unsigned int n, typename U>
729 HIP_vector_type<T, n> operator*(
730 const HIP_vector_type<T, n>& x, U y) noexcept
732 return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
734 template<typename T, unsigned int n, typename U>
738 HIP_vector_type<T, n> operator*(
739 U x, const HIP_vector_type<T, n>& y) noexcept
741 return HIP_vector_type<T, n>{x} *= y;
744 template<typename T, unsigned int n, typename U>
748 HIP_vector_type<T, n> operator/(
749 const HIP_vector_type<T, n>& x, U y) noexcept
751 return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
753 template<typename T, unsigned int n, typename U>
757 HIP_vector_type<T, n> operator/(
758 U x, const HIP_vector_type<T, n>& y) noexcept
760 return HIP_vector_type<T, n>{x} /= y;
767 bool _hip_any_zero(const V& x, int n) noexcept
770 (n == -1) ? true : ((x[n] == 0) ? false : _hip_any_zero(x, n - 1));
773 template<typename T, unsigned int n>
778 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
780 return _hip_any_zero(x.data == y.data, n - 1);
782 template<typename T, unsigned int n, typename U>
786 bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
788 return x == HIP_vector_type<T, n>{y};
790 template<typename T, unsigned int n, typename U>
794 bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
796 return HIP_vector_type<T, n>{x} == y;
799 template<typename T, unsigned int n>
804 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
808 template<typename T, unsigned int n, typename U>
812 bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
816 template<typename T, unsigned int n, typename U>
820 bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
828 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
832 HIP_vector_type<T, n> operator%(
833 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
835 return HIP_vector_type<T, n>{x} %= y;
841 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
845 HIP_vector_type<T, n> operator%(
846 const HIP_vector_type<T, n>& x, U y) noexcept
848 return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
854 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
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;
867 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
871 HIP_vector_type<T, n> operator^(
872 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
874 return HIP_vector_type<T, n>{x} ^= y;
880 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
884 HIP_vector_type<T, n> operator^(
885 const HIP_vector_type<T, n>& x, U y) noexcept
887 return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
893 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
897 HIP_vector_type<T, n> operator^(
898 U x, const HIP_vector_type<T, n>& y) noexcept
900 return HIP_vector_type<T, n>{x} ^= y;
906 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
910 HIP_vector_type<T, n> operator|(
911 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
913 return HIP_vector_type<T, n>{x} |= y;
919 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
923 HIP_vector_type<T, n> operator|(
924 const HIP_vector_type<T, n>& x, U y) noexcept
926 return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
932 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
936 HIP_vector_type<T, n> operator|(
937 U x, const HIP_vector_type<T, n>& y) noexcept
939 return HIP_vector_type<T, n>{x} |= y;
945 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
949 HIP_vector_type<T, n> operator&(
950 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
952 return HIP_vector_type<T, n>{x} &= y;
958 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
962 HIP_vector_type<T, n> operator&(
963 const HIP_vector_type<T, n>& x, U y) noexcept
965 return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
971 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
975 HIP_vector_type<T, n> operator&(
976 U x, const HIP_vector_type<T, n>& y) noexcept
978 return HIP_vector_type<T, n>{x} &= y;
984 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
988 HIP_vector_type<T, n> operator>>(
989 const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
991 return HIP_vector_type<T, n>{x} >>= y;
997 typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1001 HIP_vector_type<T, n> operator>>(
1002 const HIP_vector_type<T, n>& x, U y) noexcept
1004 return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
1010 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1014 HIP_vector_type<T, n> operator>>(
1015 U x,
const HIP_vector_type<T, n>& y)
noexcept
1017 return HIP_vector_type<T, n>{x} >>= y;
1023 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1027 HIP_vector_type<T, n> operator<<(
1028 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y)
noexcept
1030 return HIP_vector_type<T, n>{x} <<= y;
1036 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1040 HIP_vector_type<T, n> operator<<(
1041 const HIP_vector_type<T, n>& x, U y)
noexcept
1043 return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
1049 typename std::enable_if<std::is_arithmetic<U>::value>::type,
1050 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
1054 HIP_vector_type<T, n> operator<<(
1055 U x,
const HIP_vector_type<T, n>& y)
noexcept
1057 return HIP_vector_type<T, n>{x} <<= y;
1063 template <
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
1064 __forceinline__ __HOST_DEVICE__
typename std::enable_if<(rankT == 1 && rankU >= 1),
1065 const HIP_vector_type<T, rankT>>::type
1066 __hipMapVector(
const HIP_vector_type<U, rankU>& u) {
1067 return HIP_vector_type<T, rankT>(
static_cast<T
>(u.x));
1070 template <
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
1071 __forceinline__ __HOST_DEVICE__
typename std::enable_if<(rankT == 2 && rankU == 1),
1072 const HIP_vector_type<T, rankT>>::type
1073 __hipMapVector(
const HIP_vector_type<U, rankU>& u) {
1074 return HIP_vector_type<T, rankT> (
static_cast<T
>(u.x),
static_cast<T
>(0));
1077 template <
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
1078 __forceinline__ __HOST_DEVICE__
typename std::enable_if<(rankT == 2 && rankU >= 2),
1079 const HIP_vector_type<T, rankT>>::type
1080 __hipMapVector(
const HIP_vector_type<U, rankU>& u) {
1081 return HIP_vector_type<T, rankT> (
static_cast<T
>(u.x),
static_cast<T
>(u.y));
1084 template <
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
1085 __forceinline__ __HOST_DEVICE__
typename std::enable_if<(rankT == 4 && rankU == 1),
1086 const HIP_vector_type<T, rankT>>::type
1087 __hipMapVector(
const HIP_vector_type<U, rankU>& u) {
1088 return HIP_vector_type<T, rankT> (
static_cast<T
>(u.x),
static_cast<T
>(0),
1089 static_cast<T
>(0),
static_cast<T
>(0));
1092 template <
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
1093 __forceinline__ __HOST_DEVICE__
typename std::enable_if<(rankT == 4 && rankU == 2),
1094 const HIP_vector_type<T, rankT>>::type
1095 __hipMapVector(
const HIP_vector_type<U, rankU>& u) {
1096 return HIP_vector_type<T, rankT>(
static_cast<T
>(u.x),
static_cast<T
>(u.y),
1097 static_cast<T
>(0),
static_cast<T
>(0));
1100 template <
typename T,
unsigned int rankT,
typename U,
unsigned int rankU>
1101 __forceinline__ __HOST_DEVICE__
typename std::enable_if<(rankT == 4 && rankU == 4),
1102 const HIP_vector_type<T, rankT>>::type
1103 __hipMapVector(
const HIP_vector_type<U, rankU>& u) {
1104 return HIP_vector_type<T, rankT> (
static_cast<T
>(u.x),
static_cast<T
>(u.y),
1105 static_cast<T
>(u.z),
static_cast<T
>(u.w));
1108 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
1109 using CUDA_name##1 = HIP_vector_type<T, 1>;\
1110 using CUDA_name##2 = HIP_vector_type<T, 2>;\
1111 using CUDA_name##3 = HIP_vector_type<T, 3>;\
1112 using CUDA_name##4 = HIP_vector_type<T, 4>;
1114 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
1135__MAKE_VECTOR_TYPE__(uchar,
unsigned char);
1136__MAKE_VECTOR_TYPE__(
char,
char);
1137__MAKE_VECTOR_TYPE__(ushort,
unsigned short);
1138__MAKE_VECTOR_TYPE__(
short,
short);
1139__MAKE_VECTOR_TYPE__(uint,
unsigned int);
1140__MAKE_VECTOR_TYPE__(
int,
int);
1141__MAKE_VECTOR_TYPE__(ulong,
unsigned long);
1142__MAKE_VECTOR_TYPE__(
long,
long);
1143__MAKE_VECTOR_TYPE__(ulonglong,
unsigned long long);
1144__MAKE_VECTOR_TYPE__(longlong,
long long);
1145__MAKE_VECTOR_TYPE__(
float,
float);
1146__MAKE_VECTOR_TYPE__(
double,
double);
1150#if defined(_MSC_VER)
1151#include <mmintrin.h>
1152#include <xmmintrin.h>
1153#include <emmintrin.h>
1154#include <immintrin.h>
1208 unsigned char data[2];
1217 unsigned char data[4];
1225 unsigned char data[3];
1275 unsigned short data;
1282 unsigned short data[2];
1299 unsigned short data[3];
1373 unsigned int data[3];
1447 unsigned int data[3];
1658 unsigned char data[2];
1667 unsigned char data[4];
1670 unsigned char data[8];
1673 unsigned char data[16];
1681 unsigned char data[3];
1725 unsigned short data;
1732 unsigned short data[2];
1741 unsigned short data[4];
1744 unsigned short data[8];
1747 unsigned short data[16];
1755 unsigned short data[3];
1806 unsigned int data[2];
1815 unsigned int data[4];
1818 unsigned int data[8];
1821 unsigned int data[16];
1829 unsigned int data[3];
1880 unsigned long data[2];
1889 unsigned long data[4];
1892 unsigned long data[8];
1895 unsigned long data[16];
1903 unsigned long data[3];
1945 unsigned long long x;
1947 unsigned long long data;
1951 unsigned long long x;
1952 unsigned long long y;
1954 unsigned long long data[2];
1958 unsigned long long x;
1959 unsigned long long y;
1960 unsigned long long z;
1961 unsigned long long w;
1963 unsigned long long data[4];
1966 unsigned long long data[8];
1969 unsigned long long data[16];
1973 unsigned long long x;
1974 unsigned long long y;
1975 unsigned long long z;
1977 unsigned long long data[3];
2058#define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
2059 static inline __HOST_DEVICE__ type make_##type(comp x) { \
2064#define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
2065 static inline __HOST_DEVICE__ type make_##type(comp x, comp y) { \
2070#define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
2071 static inline __HOST_DEVICE__ type make_##type(comp x, comp y, comp z) { \
2076#define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
2077 static inline __HOST_DEVICE__ type make_##type(comp x, comp y, comp z, comp w) { \
2078 type r{x, y, z, w}; \
2082#define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
2083 static inline __HOST_DEVICE__ type make_##type(comp x) { \
2089#define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
2090 static inline __HOST_DEVICE__ type make_##type(comp x, comp y) { \
2097#define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
2098 static inline __HOST_DEVICE__ type make_##type(comp x, comp y, comp z) { \
2106#define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
2107 static inline __HOST_DEVICE__ type make_##type(comp x, comp y, comp z, comp w) { \
2117DECLOP_MAKE_ONE_COMPONENT(
unsigned char,
uchar1);
2118DECLOP_MAKE_TWO_COMPONENT(
unsigned char,
uchar2);
2119DECLOP_MAKE_THREE_COMPONENT(
unsigned char,
uchar3);
2120DECLOP_MAKE_FOUR_COMPONENT(
unsigned char,
uchar4);
2122DECLOP_MAKE_ONE_COMPONENT(
signed char,
char1);
2123DECLOP_MAKE_TWO_COMPONENT(
signed char,
char2);
2124DECLOP_MAKE_THREE_COMPONENT(
signed char,
char3);
2125DECLOP_MAKE_FOUR_COMPONENT(
signed char,
char4);
2127DECLOP_MAKE_ONE_COMPONENT(
unsigned short,
ushort1);
2128DECLOP_MAKE_TWO_COMPONENT(
unsigned short,
ushort2);
2129DECLOP_MAKE_THREE_COMPONENT(
unsigned short,
ushort3);
2130DECLOP_MAKE_FOUR_COMPONENT(
unsigned short,
ushort4);
2132DECLOP_MAKE_ONE_COMPONENT(
signed short,
short1);
2133DECLOP_MAKE_TWO_COMPONENT(
signed short,
short2);
2134DECLOP_MAKE_THREE_COMPONENT(
signed short,
short3);
2135DECLOP_MAKE_FOUR_COMPONENT(
signed short,
short4);
2137DECLOP_MAKE_ONE_COMPONENT(
unsigned int,
uint1);
2138DECLOP_MAKE_TWO_COMPONENT(
unsigned int,
uint2);
2139DECLOP_MAKE_THREE_COMPONENT(
unsigned int,
uint3);
2140DECLOP_MAKE_FOUR_COMPONENT(
unsigned int,
uint4);
2142DECLOP_MAKE_ONE_COMPONENT(
signed int,
int1);
2143DECLOP_MAKE_TWO_COMPONENT(
signed int,
int2);
2144DECLOP_MAKE_THREE_COMPONENT(
signed int,
int3);
2145DECLOP_MAKE_FOUR_COMPONENT(
signed int,
int4);
2147DECLOP_MAKE_ONE_COMPONENT(
float,
float1);
2148DECLOP_MAKE_TWO_COMPONENT(
float,
float2);
2149DECLOP_MAKE_THREE_COMPONENT(
float,
float3);
2150DECLOP_MAKE_FOUR_COMPONENT(
float,
float4);
2152DECLOP_MAKE_ONE_COMPONENT(
double,
double1);
2153DECLOP_MAKE_TWO_COMPONENT(
double,
double2);
2154DECLOP_MAKE_THREE_COMPONENT(
double,
double3);
2155DECLOP_MAKE_FOUR_COMPONENT(
double,
double4);
2157DECLOP_MAKE_ONE_COMPONENT(
unsigned long,
ulong1);
2158DECLOP_MAKE_TWO_COMPONENT(
unsigned long,
ulong2);
2159DECLOP_MAKE_THREE_COMPONENT(
unsigned long,
ulong3);
2160DECLOP_MAKE_FOUR_COMPONENT(
unsigned long,
ulong4);
2162DECLOP_MAKE_ONE_COMPONENT(
signed long,
long1);
2163DECLOP_MAKE_TWO_COMPONENT(
signed long,
long2);
2164DECLOP_MAKE_THREE_COMPONENT(
signed long,
long3);
2165DECLOP_MAKE_FOUR_COMPONENT(
signed long,
long4);
2167DECLOP_MAKE_ONE_COMPONENT(
unsigned long long,
ulonglong1);
2168DECLOP_MAKE_TWO_COMPONENT(
unsigned long long,
ulonglong2);
2169DECLOP_MAKE_THREE_COMPONENT(
unsigned long long,
ulonglong3);
2170DECLOP_MAKE_FOUR_COMPONENT(
unsigned long long,
ulonglong4);
2172DECLOP_MAKE_ONE_COMPONENT(
signed long long,
longlong1);
2173DECLOP_MAKE_TWO_COMPONENT(
signed long long,
longlong2);
2174DECLOP_MAKE_THREE_COMPONENT(
signed long long,
longlong3);
2175DECLOP_MAKE_FOUR_COMPONENT(
signed long long,
longlong4);
Definition amd_hip_vector_types.h:1610
Definition amd_hip_vector_types.h:1616
Definition amd_hip_vector_types.h:1623
Definition amd_hip_vector_types.h:1632
Definition amd_hip_vector_types.h:1635
Definition amd_hip_vector_types.h:1638
Definition amd_hip_vector_types.h:1647
Definition amd_hip_vector_types.h:1653
Definition amd_hip_vector_types.h:1660
Definition amd_hip_vector_types.h:1669
Definition amd_hip_vector_types.h:1672
Definition amd_hip_vector_types.h:1675
Definition amd_hip_vector_types.h:1684
Definition amd_hip_vector_types.h:1690
Definition amd_hip_vector_types.h:1697
Definition amd_hip_vector_types.h:1706
Definition amd_hip_vector_types.h:1709
Definition amd_hip_vector_types.h:1712
Definition amd_hip_vector_types.h:1721
Definition amd_hip_vector_types.h:1727
Definition amd_hip_vector_types.h:1734
Definition amd_hip_vector_types.h:1743
Definition amd_hip_vector_types.h:1746
Definition amd_hip_vector_types.h:1749
Definition amd_hip_vector_types.h:1758
Definition amd_hip_vector_types.h:1764
Definition amd_hip_vector_types.h:1771
Definition amd_hip_vector_types.h:1780
Definition amd_hip_vector_types.h:1783
Definition amd_hip_vector_types.h:1786
Definition amd_hip_vector_types.h:1795
Definition amd_hip_vector_types.h:1801
Definition amd_hip_vector_types.h:1808
Definition amd_hip_vector_types.h:1817
Definition amd_hip_vector_types.h:1820
Definition amd_hip_vector_types.h:1823
Definition amd_hip_vector_types.h:1832
Definition amd_hip_vector_types.h:1838
Definition amd_hip_vector_types.h:1845
Definition amd_hip_vector_types.h:1854
Definition amd_hip_vector_types.h:1857
Definition amd_hip_vector_types.h:1860
Definition amd_hip_vector_types.h:1869
Definition amd_hip_vector_types.h:1875
Definition amd_hip_vector_types.h:1882
Definition amd_hip_vector_types.h:1891
Definition amd_hip_vector_types.h:1894
Definition amd_hip_vector_types.h:1897
Definition amd_hip_vector_types.h:1906
Definition amd_hip_vector_types.h:1912
Definition amd_hip_vector_types.h:1919
Definition amd_hip_vector_types.h:1928
Definition amd_hip_vector_types.h:1931
Definition amd_hip_vector_types.h:1934
Definition amd_hip_vector_types.h:1943
Definition amd_hip_vector_types.h:1949
Definition amd_hip_vector_types.h:1956
Definition amd_hip_vector_types.h:1965
Definition amd_hip_vector_types.h:1968
Definition amd_hip_vector_types.h:1971
Definition amd_hip_vector_types.h:1980
Definition amd_hip_vector_types.h:1986
Definition amd_hip_vector_types.h:1993
Definition amd_hip_vector_types.h:2002
Definition amd_hip_vector_types.h:2005
Definition amd_hip_vector_types.h:2008
Definition amd_hip_vector_types.h:2017
Definition amd_hip_vector_types.h:2023
Definition amd_hip_vector_types.h:2030
Definition amd_hip_vector_types.h:2039
Definition amd_hip_vector_types.h:2042
Definition amd_hip_vector_types.h:2045