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 defined(__clang__) 39 #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) 40 #elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax. 41 #define __ROUND_UP_TO_NEXT_POT__(x) \ 42 (1 << (31 - __builtin_clz(x) + (x > (1 << (31 - __builtin_clz(x)))))) 43 #define __NATIVE_VECTOR__(n, T) \ 44 __attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T)))) 47 #if defined(__cplusplus) 48 #include <type_traits> 50 template<
typename T,
unsigned int n>
struct HIP_vector_base;
53 struct HIP_vector_base<T, 1> {
54 typedef T Native_vec_ __NATIVE_VECTOR__(1, T);
65 struct HIP_vector_base<T, 2> {
66 typedef T Native_vec_ __NATIVE_VECTOR__(2, T);
78 struct HIP_vector_base<T, 3> {
79 typedef T Native_vec_ __NATIVE_VECTOR__(3, T);
92 struct HIP_vector_base<T, 4> {
93 typedef T Native_vec_ __NATIVE_VECTOR__(4, T);
106 template<
typename T,
unsigned int rank>
107 struct HIP_vector_type :
public HIP_vector_base<T, rank> {
108 using HIP_vector_base<T, rank>::data;
109 using typename HIP_vector_base<T, rank>::Native_vec_;
112 HIP_vector_type() =
default;
115 typename std::enable_if<
116 std::is_convertible<U, T>{}>::type* =
nullptr>
118 HIP_vector_type(U x) noexcept
120 for (
auto i = 0u; i != rank; ++i) data[i] = x;
124 typename std::enable_if<
125 (rank > 1) &&
sizeof...(Us) == rank>::type* =
nullptr>
127 HIP_vector_type(Us... xs) noexcept { data = Native_vec_{
static_cast<T
>(xs)...}; }
129 HIP_vector_type(
const HIP_vector_type&) =
default;
131 HIP_vector_type(HIP_vector_type&&) =
default;
133 ~HIP_vector_type() =
default;
136 HIP_vector_type& operator=(
const HIP_vector_type&) =
default;
138 HIP_vector_type& operator=(HIP_vector_type&&) =
default;
142 HIP_vector_type& operator++() noexcept
144 return *
this += HIP_vector_type{1};
147 HIP_vector_type operator++(
int) noexcept
154 HIP_vector_type& operator--() noexcept
156 return *
this -= HIP_vector_type{1};
159 HIP_vector_type operator--(
int) noexcept
166 HIP_vector_type& operator+=(
const HIP_vector_type& x) noexcept
172 HIP_vector_type& operator-=(
const HIP_vector_type& x) noexcept
179 typename std::enable_if<
180 std::is_convertible<U, T>{}>::type* =
nullptr>
182 HIP_vector_type& operator-=(U x) noexcept
184 return *
this -= HIP_vector_type{x};
187 HIP_vector_type& operator*=(
const HIP_vector_type& x) noexcept
193 HIP_vector_type& operator/=(
const HIP_vector_type& x) noexcept
201 typename std::enable_if<std::is_signed<U>{}>::type* =
nullptr>
203 HIP_vector_type operator-() noexcept
206 tmp.data = -tmp.data;
212 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
214 HIP_vector_type operator~() noexcept
216 HIP_vector_type r{*
this};
222 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
224 HIP_vector_type& operator%=(
const HIP_vector_type& x) noexcept
231 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
233 HIP_vector_type& operator^=(
const HIP_vector_type& x) noexcept
240 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
242 HIP_vector_type& operator|=(
const HIP_vector_type& x) noexcept
249 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
251 HIP_vector_type& operator&=(
const HIP_vector_type& x) noexcept
258 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
260 HIP_vector_type& operator>>=(
const HIP_vector_type& x) noexcept
267 typename std::enable_if<std::is_integral<U>{}>::type* =
nullptr>
269 HIP_vector_type& operator<<=(
const HIP_vector_type& x) noexcept
277 template<
typename T,
unsigned int n>
279 HIP_vector_type<T, n> operator+(
280 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
282 return HIP_vector_type<T, n>{x} += y;
284 template<
typename T,
unsigned int n,
typename U>
286 HIP_vector_type<T, n> operator+(
287 const HIP_vector_type<T, n>& x, U y) noexcept
289 return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
291 template<
typename T,
unsigned int n,
typename U>
293 HIP_vector_type<T, n> operator+(
294 U x,
const HIP_vector_type<T, n>& y) noexcept
296 return HIP_vector_type<T, n>{x} += y;
299 template<
typename T,
unsigned int n>
301 HIP_vector_type<T, n> operator-(
302 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
304 return HIP_vector_type<T, n>{x} -= y;
306 template<
typename T,
unsigned int n,
typename U>
308 HIP_vector_type<T, n> operator-(
309 const HIP_vector_type<T, n>& x, U y) noexcept
311 return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
313 template<
typename T,
unsigned int n,
typename U>
315 HIP_vector_type<T, n> operator-(
316 U x,
const HIP_vector_type<T, n>& y) noexcept
318 return HIP_vector_type<T, n>{x} -= y;
321 template<
typename T,
unsigned int n>
323 HIP_vector_type<T, n> operator*(
324 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
326 return HIP_vector_type<T, n>{x} *= y;
328 template<
typename T,
unsigned int n,
typename U>
330 HIP_vector_type<T, n> operator*(
331 const HIP_vector_type<T, n>& x, U y) noexcept
333 return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
335 template<
typename T,
unsigned int n,
typename U>
337 HIP_vector_type<T, n> operator*(
338 U x,
const HIP_vector_type<T, n>& y) noexcept
340 return HIP_vector_type<T, n>{x} *= y;
343 template<
typename T,
unsigned int n>
345 HIP_vector_type<T, n> operator/(
346 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
348 return HIP_vector_type<T, n>{x} /= y;
350 template<
typename T,
unsigned int n,
typename U>
352 HIP_vector_type<T, n> operator/(
353 const HIP_vector_type<T, n>& x, U y) noexcept
355 return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
357 template<
typename T,
unsigned int n,
typename U>
359 HIP_vector_type<T, n> operator/(
360 U x,
const HIP_vector_type<T, n>& y) noexcept
362 return HIP_vector_type<T, n>{x} /= y;
365 template<
typename T,
unsigned int n>
368 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
370 auto tmp = x.data == y.data;
371 for (
auto i = 0u; i != n; ++i)
if (tmp[i] == 0)
return false;
374 template<
typename T,
unsigned int n,
typename U>
376 bool operator==(
const HIP_vector_type<T, n>& x, U y) noexcept
378 return x == HIP_vector_type<T, n>{y};
380 template<
typename T,
unsigned int n,
typename U>
382 bool operator==(U x,
const HIP_vector_type<T, n>& y) noexcept
384 return HIP_vector_type<T, n>{x} == y;
387 template<
typename T,
unsigned int n>
390 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
394 template<
typename T,
unsigned int n,
typename U>
396 bool operator!=(
const HIP_vector_type<T, n>& x, U y) noexcept
400 template<
typename T,
unsigned int n,
typename U>
402 bool operator!=(U x,
const HIP_vector_type<T, n>& y) noexcept
410 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
412 HIP_vector_type<T, n> operator%(
413 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
415 return HIP_vector_type<T, n>{x} %= y;
421 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
423 HIP_vector_type<T, n> operator%(
424 const HIP_vector_type<T, n>& x, U y) noexcept
426 return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
432 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
434 HIP_vector_type<T, n> operator%(
435 U x,
const HIP_vector_type<T, n>& y) noexcept
437 return HIP_vector_type<T, n>{x} %= y;
443 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
445 HIP_vector_type<T, n> operator^(
446 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
448 return HIP_vector_type<T, n>{x} ^= y;
454 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
456 HIP_vector_type<T, n> operator^(
457 const HIP_vector_type<T, n>& x, U y) noexcept
459 return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
465 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
467 HIP_vector_type<T, n> operator^(
468 U x,
const HIP_vector_type<T, n>& y) noexcept
470 return HIP_vector_type<T, n>{x} ^= y;
476 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
478 HIP_vector_type<T, n> operator|(
479 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
481 return HIP_vector_type<T, n>{x} |= y;
487 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
489 HIP_vector_type<T, n> operator|(
490 const HIP_vector_type<T, n>& x, U y) noexcept
492 return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
498 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
500 HIP_vector_type<T, n> operator|(
501 U x,
const HIP_vector_type<T, n>& y) noexcept
503 return HIP_vector_type<T, n>{x} |= y;
509 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
511 HIP_vector_type<T, n> operator&(
512 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
514 return HIP_vector_type<T, n>{x} &= y;
520 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
522 HIP_vector_type<T, n> operator&(
523 const HIP_vector_type<T, n>& x, U y) noexcept
525 return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
531 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
533 HIP_vector_type<T, n> operator&(
534 U x,
const HIP_vector_type<T, n>& y) noexcept
536 return HIP_vector_type<T, n>{x} &= y;
542 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
544 HIP_vector_type<T, n> operator>>(
545 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
547 return HIP_vector_type<T, n>{x} >>= y;
553 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
555 HIP_vector_type<T, n> operator>>(
556 const HIP_vector_type<T, n>& x, U y) noexcept
558 return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
564 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
566 HIP_vector_type<T, n> operator>>(
567 U x,
const HIP_vector_type<T, n>& y) noexcept
569 return HIP_vector_type<T, n>{x} >>= y;
575 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
577 HIP_vector_type<T, n> operator<<(
578 const HIP_vector_type<T, n>& x,
const HIP_vector_type<T, n>& y) noexcept
580 return HIP_vector_type<T, n>{x} <<= y;
586 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
588 HIP_vector_type<T, n> operator<<(
589 const HIP_vector_type<T, n>& x, U y) noexcept
591 return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
597 typename std::enable_if<std::is_integral<T>{}>* =
nullptr>
599 HIP_vector_type<T, n> operator<<(
600 U x,
const HIP_vector_type<T, n>& y) noexcept
602 return HIP_vector_type<T, n>{x} <<= y;
605 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 606 using CUDA_name##1 = HIP_vector_type<T, 1>;\ 607 using CUDA_name##2 = HIP_vector_type<T, 2>;\ 608 using CUDA_name##3 = HIP_vector_type<T, 3>;\ 609 using CUDA_name##4 = HIP_vector_type<T, 4>; 611 #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \ 612 typedef T CUDA_name##_impl1 __NATIVE_VECTOR__(1, T);\ 613 typedef T CUDA_name##_impl2 __NATIVE_VECTOR__(2, T);\ 614 typedef T CUDA_name##_impl3 __NATIVE_VECTOR__(3, T);\ 615 typedef T CUDA_name##_impl4 __NATIVE_VECTOR__(4, T);\ 618 CUDA_name##_impl1 data;\ 626 CUDA_name##_impl2 data;\ 635 CUDA_name##_impl3 data;\ 645 CUDA_name##_impl4 data;\ 656 __MAKE_VECTOR_TYPE__(uchar,
unsigned char);
657 __MAKE_VECTOR_TYPE__(
char,
char);
658 __MAKE_VECTOR_TYPE__(ushort,
unsigned short);
659 __MAKE_VECTOR_TYPE__(
short,
short);
660 __MAKE_VECTOR_TYPE__(uint,
unsigned int);
661 __MAKE_VECTOR_TYPE__(
int,
int);
662 __MAKE_VECTOR_TYPE__(ulong,
unsigned long);
663 __MAKE_VECTOR_TYPE__(
long,
long);
664 __MAKE_VECTOR_TYPE__(ulonglong,
unsigned long long);
665 __MAKE_VECTOR_TYPE__(longlong,
long long);
666 __MAKE_VECTOR_TYPE__(
float,
float);
667 __MAKE_VECTOR_TYPE__(
double,
double);
669 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ 670 static inline __device__ __host__ \ 671 type make_##type(comp x) { type r{x}; return r; } 673 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ 674 static inline __device__ __host__ \ 675 type make_##type(comp x, comp y) { type r{x, y}; return r; } 677 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ 678 static inline __device__ __host__ \ 679 type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; } 681 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ 682 static inline __device__ __host__ \ 683 type make_##type(comp x, comp y, comp z, comp w) { \ 684 type r{x, y, z, w}; \ 688 DECLOP_MAKE_ONE_COMPONENT(
unsigned char, uchar1);
689 DECLOP_MAKE_TWO_COMPONENT(
unsigned char, uchar2);
690 DECLOP_MAKE_THREE_COMPONENT(
unsigned char, uchar3);
691 DECLOP_MAKE_FOUR_COMPONENT(
unsigned char, uchar4);
693 DECLOP_MAKE_ONE_COMPONENT(
signed char, char1);
694 DECLOP_MAKE_TWO_COMPONENT(
signed char, char2);
695 DECLOP_MAKE_THREE_COMPONENT(
signed char, char3);
696 DECLOP_MAKE_FOUR_COMPONENT(
signed char, char4);
698 DECLOP_MAKE_ONE_COMPONENT(
unsigned short, ushort1);
699 DECLOP_MAKE_TWO_COMPONENT(
unsigned short, ushort2);
700 DECLOP_MAKE_THREE_COMPONENT(
unsigned short, ushort3);
701 DECLOP_MAKE_FOUR_COMPONENT(
unsigned short, ushort4);
703 DECLOP_MAKE_ONE_COMPONENT(
signed short, short1);
704 DECLOP_MAKE_TWO_COMPONENT(
signed short, short2);
705 DECLOP_MAKE_THREE_COMPONENT(
signed short, short3);
706 DECLOP_MAKE_FOUR_COMPONENT(
signed short, short4);
708 DECLOP_MAKE_ONE_COMPONENT(
unsigned int, uint1);
709 DECLOP_MAKE_TWO_COMPONENT(
unsigned int, uint2);
710 DECLOP_MAKE_THREE_COMPONENT(
unsigned int, uint3);
711 DECLOP_MAKE_FOUR_COMPONENT(
unsigned int, uint4);
713 DECLOP_MAKE_ONE_COMPONENT(
signed int, int1);
714 DECLOP_MAKE_TWO_COMPONENT(
signed int, int2);
715 DECLOP_MAKE_THREE_COMPONENT(
signed int, int3);
716 DECLOP_MAKE_FOUR_COMPONENT(
signed int, int4);
718 DECLOP_MAKE_ONE_COMPONENT(
float, float1);
719 DECLOP_MAKE_TWO_COMPONENT(
float, float2);
720 DECLOP_MAKE_THREE_COMPONENT(
float, float3);
721 DECLOP_MAKE_FOUR_COMPONENT(
float, float4);
723 DECLOP_MAKE_ONE_COMPONENT(
double, double1);
724 DECLOP_MAKE_TWO_COMPONENT(
double, double2);
725 DECLOP_MAKE_THREE_COMPONENT(
double, double3);
726 DECLOP_MAKE_FOUR_COMPONENT(
double, double4);
728 DECLOP_MAKE_ONE_COMPONENT(
unsigned long, ulong1);
729 DECLOP_MAKE_TWO_COMPONENT(
unsigned long, ulong2);
730 DECLOP_MAKE_THREE_COMPONENT(
unsigned long, ulong3);
731 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long, ulong4);
733 DECLOP_MAKE_ONE_COMPONENT(
signed long, long1);
734 DECLOP_MAKE_TWO_COMPONENT(
signed long, long2);
735 DECLOP_MAKE_THREE_COMPONENT(
signed long, long3);
736 DECLOP_MAKE_FOUR_COMPONENT(
signed long, long4);
738 DECLOP_MAKE_ONE_COMPONENT(
unsigned long long, ulonglong1);
739 DECLOP_MAKE_TWO_COMPONENT(
unsigned long long, ulonglong2);
740 DECLOP_MAKE_THREE_COMPONENT(
unsigned long long, ulonglong3);
741 DECLOP_MAKE_FOUR_COMPONENT(
unsigned long long, ulonglong4);
743 DECLOP_MAKE_ONE_COMPONENT(
signed long long, longlong1);
744 DECLOP_MAKE_TWO_COMPONENT(
signed long long, longlong2);
745 DECLOP_MAKE_THREE_COMPONENT(
signed long long, longlong3);
746 DECLOP_MAKE_FOUR_COMPONENT(
signed long long, longlong4);
747 #else // defined(_MSC_VER) 748 #include <mmintrin.h> 749 #include <xmmintrin.h> 750 #include <emmintrin.h> 751 #include <immintrin.h> 753 typedef union {
char data; } char1;
754 typedef union {
char data[2]; } char2;
755 typedef union {
char data[4]; } char4;
756 typedef union { char4 data; } char3;
757 typedef union { __m64 data; } char8;
758 typedef union { __m128i data; } char16;
760 typedef union {
unsigned char data; } uchar1;
761 typedef union {
unsigned char data[2]; } uchar2;
762 typedef union {
unsigned char data[4]; } uchar4;
763 typedef union { uchar4 data; } uchar3;
764 typedef union { __m64 data; } uchar8;
765 typedef union { __m128i data; } uchar16;
767 typedef union {
short data; } short1;
768 typedef union {
short data[2]; } short2;
769 typedef union { __m64 data; } short4;
770 typedef union { short4 data; } short3;
771 typedef union { __m128i data; } short8;
772 typedef union { __m128i data[2]; } short16;
774 typedef union {
unsigned short data; } ushort1;
775 typedef union {
unsigned short data[2]; } ushort2;
776 typedef union { __m64 data; } ushort4;
777 typedef union { ushort4 data; } ushort3;
778 typedef union { __m128i data; } ushort8;
779 typedef union { __m128i data[2]; } ushort16;
781 typedef union {
int data; } int1;
782 typedef union { __m64 data; } int2;
783 typedef union { __m128i data; } int4;
784 typedef union { int4 data; } int3;
785 typedef union { __m128i data[2]; } int8;
786 typedef union { __m128i data[4];} int16;
788 typedef union {
unsigned int data; } uint1;
789 typedef union { __m64 data; } uint2;
790 typedef union { __m128i data; } uint4;
791 typedef union { uint4 data; } uint3;
792 typedef union { __m128i data[2]; } uint8;
793 typedef union { __m128i data[4]; } uint16;
796 typedef union {
int data; } long1;
797 typedef union { __m64 data; } long2;
798 typedef union { __m128i data; } long4;
799 typedef union { long4 data; } long3;
800 typedef union { __m128i data[2]; } long8;
801 typedef union { __m128i data[4]; } long16;
803 typedef union {
unsigned int data; } ulong1;
804 typedef union { __m64 data; } ulong2;
805 typedef union { __m128i data; } ulong4;
806 typedef union { ulong4 data; } ulong3;
807 typedef union { __m128i data[2]; } ulong8;
808 typedef union { __m128i data[4]; } ulong16;
809 #else // defined(_WIN64) 810 typedef union { __m64 data; } long1;
811 typedef union { __m128i data; } long2;
812 typedef union { __m128i data[2]; } long4;
813 typedef union { long4 data; } long3;
814 typedef union { __m128i data[4]; } long8;
815 typedef union { __m128i data[8]; } long16;
817 typedef union { __m64 data; } ulong1;
818 typedef union { __m128i data; } ulong2;
819 typedef union { __m128i data[2]; } ulong4;
820 typedef union { ulong4 data; } ulong3;
821 typedef union { __m128i data[4]; } ulong8;
822 typedef union { __m128i data[8]; } ulong16;
823 #endif // defined(_WIN64) 825 typedef union { __m64 data; } longlong1;
826 typedef union { __m128i data; } longlong2;
827 typedef union { __m128i data[2]; } longlong4;
828 typedef union { longlong4 data; } longlong3;
829 typedef union { __m128i data[4]; } longlong8;
830 typedef union { __m128i data[8]; } longlong16;
832 typedef union { __m64 data; } ulonglong1;
833 typedef union { __m128i data; } ulonglong2;
834 typedef union { __m128i data[2]; } ulonglong4;
835 typedef union { ulonglong4 data; } ulonglong3;
836 typedef union { __m128i data[4]; } ulonglong8;
837 typedef union { __m128i data[8]; } ulonglong16;
839 typedef union {
float data; } float1;
840 typedef union { __m64 data; } float2;
841 typedef union { __m128 data; } float4;
842 typedef union { float4 data; } float3;
843 typedef union { __m256 data; } float8;
844 typedef union { __m256 data[2]; } float16;
846 typedef union {
double data; } double1;
847 typedef union { __m128d data; } double2;
848 typedef union { __m256d data; } double4;
849 typedef union { double4 data; } double3;
850 typedef union { __m256d data[2]; } double8;
851 typedef union { __m256d data[4]; } double16;
853 #endif // defined(_MSC_VER)
#define __host__
Definition: host_defines.h:41