HIP: Heterogenous-computing Interface for Portability
hip_vector_types.h
Go to the documentation of this file.
1 /*
2 Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
28 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H
29 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_VECTOR_TYPES_H
30 
31 #if defined(__HCC__) && (__hcc_workweek__ < 16032)
32 #error("This version of HIP requires a newer version of HCC.");
33 #endif
34 
36 
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))))
45 #endif
46 
47 #if defined(__cplusplus)
48  #include <type_traits>
49 
50  template<typename T, unsigned int n> struct HIP_vector_base;
51 
52  template<typename T>
53  struct HIP_vector_base<T, 1> {
54  typedef T Native_vec_ __NATIVE_VECTOR__(1, T);
55 
56  union {
57  Native_vec_ data;
58  struct {
59  T x;
60  };
61  };
62  };
63 
64  template<typename T>
65  struct HIP_vector_base<T, 2> {
66  typedef T Native_vec_ __NATIVE_VECTOR__(2, T);
67 
68  union {
69  Native_vec_ data;
70  struct {
71  T x;
72  T y;
73  };
74  };
75  };
76 
77  template<typename T>
78  struct HIP_vector_base<T, 3> {
79  typedef T Native_vec_ __NATIVE_VECTOR__(3, T);
80 
81  union {
82  Native_vec_ data;
83  struct {
84  T x;
85  T y;
86  T z;
87  };
88  };
89  };
90 
91  template<typename T>
92  struct HIP_vector_base<T, 4> {
93  typedef T Native_vec_ __NATIVE_VECTOR__(4, T);
94 
95  union {
96  Native_vec_ data;
97  struct {
98  T x;
99  T y;
100  T z;
101  T w;
102  };
103  };
104  };
105 
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_;
110 
111  __host__ __device__
112  HIP_vector_type() = default;
113  template<
114  typename U,
115  typename std::enable_if<
116  std::is_convertible<U, T>{}>::type* = nullptr>
117  __host__ __device__
118  HIP_vector_type(U x) noexcept
119  {
120  for (auto i = 0u; i != rank; ++i) data[i] = x;
121  }
122  template< // TODO: constrain based on type as well.
123  typename... Us,
124  typename std::enable_if<
125  (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
126  __host__ __device__
127  HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast<T>(xs)...}; }
128  __host__ __device__
129  HIP_vector_type(const HIP_vector_type&) = default;
130  __host__ __device__
131  HIP_vector_type(HIP_vector_type&&) = default;
132  __host__ __device__
133  ~HIP_vector_type() = default;
134 
135  __host__ __device__
136  HIP_vector_type& operator=(const HIP_vector_type&) = default;
137  __host__ __device__
138  HIP_vector_type& operator=(HIP_vector_type&&) = default;
139 
140  // Operators
141  __host__ __device__
142  HIP_vector_type& operator++() noexcept
143  {
144  return *this += HIP_vector_type{1};
145  }
146  __host__ __device__
147  HIP_vector_type operator++(int) noexcept
148  {
149  auto tmp(*this);
150  ++*this;
151  return tmp;
152  }
153  __host__ __device__
154  HIP_vector_type& operator--() noexcept
155  {
156  return *this -= HIP_vector_type{1};
157  }
158  __host__ __device__
159  HIP_vector_type operator--(int) noexcept
160  {
161  auto tmp(*this);
162  --*this;
163  return tmp;
164  }
165  __host__ __device__
166  HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
167  {
168  data += x.data;
169  return *this;
170  }
171  __host__ __device__
172  HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
173  {
174  data -= x.data;
175  return *this;
176  }
177  template<
178  typename U,
179  typename std::enable_if<
180  std::is_convertible<U, T>{}>::type* = nullptr>
181  __host__ __device__
182  HIP_vector_type& operator-=(U x) noexcept
183  {
184  return *this -= HIP_vector_type{x};
185  }
186  __host__ __device__
187  HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
188  {
189  data *= x.data;
190  return *this;
191  }
192  __host__ __device__
193  HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
194  {
195  data /= x.data;
196  return *this;
197  }
198 
199  template<
200  typename U = T,
201  typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
202  __host__ __device__
203  HIP_vector_type operator-() noexcept
204  {
205  auto tmp(*this);
206  tmp.data = -tmp.data;
207  return tmp;
208  }
209 
210  template<
211  typename U = T,
212  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
213  __host__ __device__
214  HIP_vector_type operator~() noexcept
215  {
216  HIP_vector_type r{*this};
217  r.data = ~r.data;
218  return r;
219  }
220  template<
221  typename U = T,
222  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
223  __host__ __device__
224  HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
225  {
226  data %= x.data;
227  return *this;
228  }
229  template<
230  typename U = T,
231  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
232  __host__ __device__
233  HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
234  {
235  data ^= x.data;
236  return *this;
237  }
238  template<
239  typename U = T,
240  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
241  __host__ __device__
242  HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
243  {
244  data |= x.data;
245  return *this;
246  }
247  template<
248  typename U = T,
249  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
250  __host__ __device__
251  HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
252  {
253  data &= x.data;
254  return *this;
255  }
256  template<
257  typename U = T,
258  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
259  __host__ __device__
260  HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
261  {
262  data >>= x.data;
263  return *this;
264  }
265  template<
266  typename U = T,
267  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
268  __host__ __device__
269  HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
270  {
271  data <<= x.data;
272  return *this;
273  }
274  };
275 
276 
277  template<typename T, unsigned int n>
278  __host__ __device__
279  inline
280  HIP_vector_type<T, n> operator+(
281  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
282  {
283  return HIP_vector_type<T, n>{x} += y;
284  }
285  template<typename T, unsigned int n, typename U>
286  __host__ __device__
287  inline
288  HIP_vector_type<T, n> operator+(
289  const HIP_vector_type<T, n>& x, U y) noexcept
290  {
291  return HIP_vector_type<T, n>{x} += y;
292  }
293  template<typename T, unsigned int n, typename U>
294  __host__ __device__
295  inline
296  HIP_vector_type<T, n> operator+(
297  U x, const HIP_vector_type<T, n>& y) noexcept
298  {
299  return y + x;
300  }
301 
302  template<typename T, unsigned int n>
303  __host__ __device__
304  inline
305  HIP_vector_type<T, n> operator-(
306  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
307  {
308  return HIP_vector_type<T, n>{x} -= y;
309  }
310  template<typename T, unsigned int n, typename U>
311  __host__ __device__
312  inline
313  HIP_vector_type<T, n> operator-(
314  const HIP_vector_type<T, n>& x, U y) noexcept
315  {
316  return HIP_vector_type<T, n>{x} -= y;
317  }
318  template<typename T, unsigned int n, typename U>
319  __host__ __device__
320  inline
321  HIP_vector_type<T, n> operator-(
322  U x, const HIP_vector_type<T, n>& y) noexcept
323  {
324  return HIP_vector_type<T, n>{x} -= y;
325  }
326 
327  template<typename T, unsigned int n>
328  __host__ __device__
329  inline
330  HIP_vector_type<T, n> operator*(
331  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
332  {
333  return HIP_vector_type<T, n>{x} *= y;
334  }
335  template<typename T, unsigned int n, typename U>
336  __host__ __device__
337  inline
338  HIP_vector_type<T, n> operator*(
339  const HIP_vector_type<T, n>& x, U y) noexcept
340  {
341  return HIP_vector_type<T, n>{x} *= y;
342  }
343  template<typename T, unsigned int n, typename U>
344  __host__ __device__
345  inline
346  HIP_vector_type<T, n> operator*(
347  U x, const HIP_vector_type<T, n>& y) noexcept
348  {
349  return y * x;
350  }
351 
352  template<typename T, unsigned int n>
353  __host__ __device__
354  inline
355  HIP_vector_type<T, n> operator/(
356  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
357  {
358  return HIP_vector_type<T, n>{x} /= y;
359  }
360  template<typename T, unsigned int n, typename U>
361  __host__ __device__
362  inline
363  HIP_vector_type<T, n> operator/(
364  const HIP_vector_type<T, n>& x, U y) noexcept
365  {
366  return HIP_vector_type<T, n>{x} /= y;
367  }
368  template<typename T, unsigned int n, typename U>
369  __host__ __device__
370  inline
371  HIP_vector_type<T, n> operator/(
372  U x, const HIP_vector_type<T, n>& y) noexcept
373  {
374  return HIP_vector_type<T, n>{x} /= y;
375  }
376 
377  template<typename T, unsigned int n>
378  __host__ __device__
379  inline
380  bool operator==(
381  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
382  {
383  auto tmp = x.data == y.data;
384  for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false;
385  return true;
386  }
387  template<typename T, unsigned int n, typename U>
388  __host__ __device__
389  inline
390  bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
391  {
392  return x == HIP_vector_type<T, n>{y};
393  }
394  template<typename T, unsigned int n, typename U>
395  __host__ __device__
396  inline
397  bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
398  {
399  return HIP_vector_type<T, n>{x} == y;
400  }
401 
402  template<typename T, unsigned int n>
403  __host__ __device__
404  inline
405  bool operator!=(
406  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
407  {
408  return !(x == y);
409  }
410  template<typename T, unsigned int n, typename U>
411  __host__ __device__
412  inline
413  bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
414  {
415  return !(x == y);
416  }
417  template<typename T, unsigned int n, typename U>
418  __host__ __device__
419  inline
420  bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
421  {
422  return !(x == y);
423  }
424 
425  template<
426  typename T,
427  unsigned int n,
428  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
429  inline
430  HIP_vector_type<T, n> operator%(
431  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
432  {
433  return HIP_vector_type<T, n>{x} %= y;
434  }
435  template<
436  typename T,
437  unsigned int n,
438  typename U,
439  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
440  inline
441  HIP_vector_type<T, n> operator%(
442  const HIP_vector_type<T, n>& x, U y) noexcept
443  {
444  return HIP_vector_type<T, n>{x} %= y;
445  }
446  template<
447  typename T,
448  unsigned int n,
449  typename U,
450  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
451  inline
452  HIP_vector_type<T, n> operator%(
453  U x, const HIP_vector_type<T, n>& y) noexcept
454  {
455  return HIP_vector_type<T, n>{x} %= y;
456  }
457 
458  template<
459  typename T,
460  unsigned int n,
461  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
462  inline
463  HIP_vector_type<T, n> operator^(
464  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
465  {
466  return HIP_vector_type<T, n>{x} ^= y;
467  }
468  template<
469  typename T,
470  unsigned int n,
471  typename U,
472  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
473  inline
474  HIP_vector_type<T, n> operator^(
475  const HIP_vector_type<T, n>& x, U y) noexcept
476  {
477  return HIP_vector_type<T, n>{x} ^= y;
478  }
479  template<
480  typename T,
481  unsigned int n,
482  typename U,
483  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
484  inline
485  HIP_vector_type<T, n> operator^(
486  U x, const HIP_vector_type<T, n>& y) noexcept
487  {
488  return HIP_vector_type<T, n>{x} ^= y;
489  }
490 
491  template<
492  typename T,
493  unsigned int n,
494  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
495  inline
496  HIP_vector_type<T, n> operator|(
497  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
498  {
499  return HIP_vector_type<T, n>{x} |= y;
500  }
501  template<
502  typename T,
503  unsigned int n,
504  typename U,
505  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
506  inline
507  HIP_vector_type<T, n> operator|(
508  const HIP_vector_type<T, n>& x, U y) noexcept
509  {
510  return HIP_vector_type<T, n>{x} |= y;
511  }
512  template<
513  typename T,
514  unsigned int n,
515  typename U,
516  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
517  inline
518  HIP_vector_type<T, n> operator|(
519  U x, const HIP_vector_type<T, n>& y) noexcept
520  {
521  return HIP_vector_type<T, n>{x} |= y;
522  }
523 
524  template<
525  typename T,
526  unsigned int n,
527  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
528  inline
529  HIP_vector_type<T, n> operator&(
530  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
531  {
532  return HIP_vector_type<T, n>{x} &= y;
533  }
534  template<
535  typename T,
536  unsigned int n,
537  typename U,
538  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
539  inline
540  HIP_vector_type<T, n> operator&(
541  const HIP_vector_type<T, n>& x, U y) noexcept
542  {
543  return HIP_vector_type<T, n>{x} &= y;
544  }
545  template<
546  typename T,
547  unsigned int n,
548  typename U,
549  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
550  inline
551  HIP_vector_type<T, n> operator&(
552  U x, const HIP_vector_type<T, n>& y) noexcept
553  {
554  return HIP_vector_type<T, n>{x} &= y;
555  }
556 
557  template<
558  typename T,
559  unsigned int n,
560  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
561  inline
562  HIP_vector_type<T, n> operator>>(
563  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
564  {
565  return HIP_vector_type<T, n>{x} >>= y;
566  }
567  template<
568  typename T,
569  unsigned int n,
570  typename U,
571  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
572  inline
573  HIP_vector_type<T, n> operator>>(
574  const HIP_vector_type<T, n>& x, U y) noexcept
575  {
576  return HIP_vector_type<T, n>{x} >>= y;
577  }
578  template<
579  typename T,
580  unsigned int n,
581  typename U,
582  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
583  inline
584  HIP_vector_type<T, n> operator>>(
585  U x, const HIP_vector_type<T, n>& y) noexcept
586  {
587  return HIP_vector_type<T, n>{x} >>= y;
588  }
589 
590  template<
591  typename T,
592  unsigned int n,
593  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
594  inline
595  HIP_vector_type<T, n> operator<<(
596  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
597  {
598  return HIP_vector_type<T, n>{x} <<= y;
599  }
600  template<
601  typename T,
602  unsigned int n,
603  typename U,
604  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
605  inline
606  HIP_vector_type<T, n> operator<<(
607  const HIP_vector_type<T, n>& x, U y) noexcept
608  {
609  return HIP_vector_type<T, n>{x} <<= y;
610  }
611  template<
612  typename T,
613  unsigned int n,
614  typename U,
615  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
616  inline
617  HIP_vector_type<T, n> operator<<(
618  U x, const HIP_vector_type<T, n>& y) noexcept
619  {
620  return HIP_vector_type<T, n>{x} <<= y;
621  }
622 
623  #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
624  using CUDA_name##1 = HIP_vector_type<T, 1>;\
625  using CUDA_name##2 = HIP_vector_type<T, 2>;\
626  using CUDA_name##3 = HIP_vector_type<T, 3>;\
627  using CUDA_name##4 = HIP_vector_type<T, 4>;
628 #else
629  #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
630  typedef T CUDA_name##_impl1 __NATIVE_VECTOR__(1, T);\
631  typedef T CUDA_name##_impl2 __NATIVE_VECTOR__(2, T);\
632  typedef T CUDA_name##_impl3 __NATIVE_VECTOR__(3, T);\
633  typedef T CUDA_name##_impl4 __NATIVE_VECTOR__(4, T);\
634  typedef struct {\
635  union {\
636  CUDA_name##_impl1 data;\
637  struct {\
638  T x;\
639  };\
640  };\
641  } CUDA_name##1;\
642  typedef struct {\
643  union {\
644  CUDA_name##_impl2 data;\
645  struct {\
646  T x;\
647  T y;\
648  };\
649  };\
650  } CUDA_name##2;\
651  typedef struct {\
652  union {\
653  CUDA_name##_impl3 data;\
654  struct {\
655  T x;\
656  T y;\
657  T z;\
658  };\
659  };\
660  } CUDA_name##3;\
661  typedef struct {\
662  union {\
663  CUDA_name##_impl4 data;\
664  struct {\
665  T x;\
666  T y;\
667  T z;\
668  T w;\
669  };\
670  };\
671  } CUDA_name##4;
672 #endif
673 
674 __MAKE_VECTOR_TYPE__(uchar, unsigned char);
675 __MAKE_VECTOR_TYPE__(char, char);
676 __MAKE_VECTOR_TYPE__(ushort, unsigned short);
677 __MAKE_VECTOR_TYPE__(short, short);
678 __MAKE_VECTOR_TYPE__(uint, unsigned int);
679 __MAKE_VECTOR_TYPE__(int, int);
680 __MAKE_VECTOR_TYPE__(ulong, unsigned long);
681 __MAKE_VECTOR_TYPE__(long, long);
682 __MAKE_VECTOR_TYPE__(ulonglong, unsigned long long);
683 __MAKE_VECTOR_TYPE__(longlong, long long);
684 __MAKE_VECTOR_TYPE__(float, float);
685 __MAKE_VECTOR_TYPE__(double, double);
686 
687 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
688  __device__ __host__ \
689  static \
690  inline \
691  type make_##type(comp x) { type r = {x}; return r; }
692 
693 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
694  __device__ __host__ \
695  static \
696  inline \
697  type make_##type(comp x, comp y) { type r = {x, y}; return r; }
698 
699 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
700  __device__ __host__ \
701  static \
702  inline \
703  type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; }
704 
705 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
706  __device__ __host__ \
707  static \
708  inline \
709  type make_##type(comp x, comp y, comp z, comp w) { \
710  type r = {x, y, z, w}; \
711  return r; \
712  }
713 
714 DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1);
715 DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2);
716 DECLOP_MAKE_THREE_COMPONENT(unsigned char, uchar3);
717 DECLOP_MAKE_FOUR_COMPONENT(unsigned char, uchar4);
718 
719 DECLOP_MAKE_ONE_COMPONENT(signed char, char1);
720 DECLOP_MAKE_TWO_COMPONENT(signed char, char2);
721 DECLOP_MAKE_THREE_COMPONENT(signed char, char3);
722 DECLOP_MAKE_FOUR_COMPONENT(signed char, char4);
723 
724 DECLOP_MAKE_ONE_COMPONENT(unsigned short, ushort1);
725 DECLOP_MAKE_TWO_COMPONENT(unsigned short, ushort2);
726 DECLOP_MAKE_THREE_COMPONENT(unsigned short, ushort3);
727 DECLOP_MAKE_FOUR_COMPONENT(unsigned short, ushort4);
728 
729 DECLOP_MAKE_ONE_COMPONENT(signed short, short1);
730 DECLOP_MAKE_TWO_COMPONENT(signed short, short2);
731 DECLOP_MAKE_THREE_COMPONENT(signed short, short3);
732 DECLOP_MAKE_FOUR_COMPONENT(signed short, short4);
733 
734 DECLOP_MAKE_ONE_COMPONENT(unsigned int, uint1);
735 DECLOP_MAKE_TWO_COMPONENT(unsigned int, uint2);
736 DECLOP_MAKE_THREE_COMPONENT(unsigned int, uint3);
737 DECLOP_MAKE_FOUR_COMPONENT(unsigned int, uint4);
738 
739 DECLOP_MAKE_ONE_COMPONENT(signed int, int1);
740 DECLOP_MAKE_TWO_COMPONENT(signed int, int2);
741 DECLOP_MAKE_THREE_COMPONENT(signed int, int3);
742 DECLOP_MAKE_FOUR_COMPONENT(signed int, int4);
743 
744 DECLOP_MAKE_ONE_COMPONENT(float, float1);
745 DECLOP_MAKE_TWO_COMPONENT(float, float2);
746 DECLOP_MAKE_THREE_COMPONENT(float, float3);
747 DECLOP_MAKE_FOUR_COMPONENT(float, float4);
748 
749 DECLOP_MAKE_ONE_COMPONENT(double, double1);
750 DECLOP_MAKE_TWO_COMPONENT(double, double2);
751 DECLOP_MAKE_THREE_COMPONENT(double, double3);
752 DECLOP_MAKE_FOUR_COMPONENT(double, double4);
753 
754 DECLOP_MAKE_ONE_COMPONENT(unsigned long, ulong1);
755 DECLOP_MAKE_TWO_COMPONENT(unsigned long, ulong2);
756 DECLOP_MAKE_THREE_COMPONENT(unsigned long, ulong3);
757 DECLOP_MAKE_FOUR_COMPONENT(unsigned long, ulong4);
758 
759 DECLOP_MAKE_ONE_COMPONENT(signed long, long1);
760 DECLOP_MAKE_TWO_COMPONENT(signed long, long2);
761 DECLOP_MAKE_THREE_COMPONENT(signed long, long3);
762 DECLOP_MAKE_FOUR_COMPONENT(signed long, long4);
763 
764 DECLOP_MAKE_ONE_COMPONENT(unsigned long long, ulonglong1);
765 DECLOP_MAKE_TWO_COMPONENT(unsigned long long, ulonglong2);
766 DECLOP_MAKE_THREE_COMPONENT(unsigned long long, ulonglong3);
767 DECLOP_MAKE_FOUR_COMPONENT(unsigned long long, ulonglong4);
768 
769 DECLOP_MAKE_ONE_COMPONENT(signed long long, longlong1);
770 DECLOP_MAKE_TWO_COMPONENT(signed long long, longlong2);
771 DECLOP_MAKE_THREE_COMPONENT(signed long long, longlong3);
772 DECLOP_MAKE_FOUR_COMPONENT(signed long long, longlong4);
773 #else // defined(_MSC_VER)
774 #include <mmintrin.h>
775 #include <xmmintrin.h>
776 #include <emmintrin.h>
777 #include <immintrin.h>
778 
779 typedef union { char data; } char1;
780 typedef union { char data[2]; } char2;
781 typedef union { char data[4]; } char4;
782 typedef union { char4 data; } char3;
783 typedef union { __m64 data; } char8;
784 typedef union { __m128i data; } char16;
785 
786 typedef union { unsigned char data; } uchar1;
787 typedef union { unsigned char data[2]; } uchar2;
788 typedef union { unsigned char data[4]; } uchar4;
789 typedef union { uchar4 data; } uchar3;
790 typedef union { __m64 data; } uchar8;
791 typedef union { __m128i data; } uchar16;
792 
793 typedef union { short data; } short1;
794 typedef union { short data[2]; } short2;
795 typedef union { __m64 data; } short4;
796 typedef union { short4 data; } short3;
797 typedef union { __m128i data; } short8;
798 typedef union { __m128i data[2]; } short16;
799 
800 typedef union { unsigned short data; } ushort1;
801 typedef union { unsigned short data[2]; } ushort2;
802 typedef union { __m64 data; } ushort4;
803 typedef union { ushort4 data; } ushort3;
804 typedef union { __m128i data; } ushort8;
805 typedef union { __m128i data[2]; } ushort16;
806 
807 typedef union { int data; } int1;
808 typedef union { __m64 data; } int2;
809 typedef union { __m128i data; } int4;
810 typedef union { int4 data; } int3;
811 typedef union { __m128i data[2]; } int8;
812 typedef union { __m128i data[4];} int16;
813 
814 typedef union { unsigned int data; } uint1;
815 typedef union { __m64 data; } uint2;
816 typedef union { __m128i data; } uint4;
817 typedef union { uint4 data; } uint3;
818 typedef union { __m128i data[2]; } uint8;
819 typedef union { __m128i data[4]; } uint16;
820 
821 #if !defined(_WIN64)
822 typedef union { int data; } long1;
823 typedef union { __m64 data; } long2;
824 typedef union { __m128i data; } long4;
825 typedef union { long4 data; } long3;
826 typedef union { __m128i data[2]; } long8;
827 typedef union { __m128i data[4]; } long16;
828 
829 typedef union { unsigned int data; } ulong1;
830 typedef union { __m64 data; } ulong2;
831 typedef union { __m128i data; } ulong4;
832 typedef union { ulong4 data; } ulong3;
833 typedef union { __m128i data[2]; } ulong8;
834 typedef union { __m128i data[4]; } ulong16;
835 #else // defined(_WIN64)
836 typedef union { __m64 data; } long1;
837 typedef union { __m128i data; } long2;
838 typedef union { __m128i data[2]; } long4;
839 typedef union { long4 data; } long3;
840 typedef union { __m128i data[4]; } long8;
841 typedef union { __m128i data[8]; } long16;
842 
843 typedef union { __m64 data; } ulong1;
844 typedef union { __m128i data; } ulong2;
845 typedef union { __m128i data[2]; } ulong4;
846 typedef union { ulong4 data; } ulong3;
847 typedef union { __m128i data[4]; } ulong8;
848 typedef union { __m128i data[8]; } ulong16;
849 #endif // defined(_WIN64)
850 
851 typedef union { __m64 data; } longlong1;
852 typedef union { __m128i data; } longlong2;
853 typedef union { __m128i data[2]; } longlong4;
854 typedef union { longlong4 data; } longlong3;
855 typedef union { __m128i data[4]; } longlong8;
856 typedef union { __m128i data[8]; } longlong16;
857 
858 typedef union { __m64 data; } ulonglong1;
859 typedef union { __m128i data; } ulonglong2;
860 typedef union { __m128i data[2]; } ulonglong4;
861 typedef union { ulonglong4 data; } ulonglong3;
862 typedef union { __m128i data[4]; } ulonglong8;
863 typedef union { __m128i data[8]; } ulonglong16;
864 
865 typedef union { float data; } float1;
866 typedef union { __m64 data; } float2;
867 typedef union { __m128 data; } float4;
868 typedef union { float4 data; } float3;
869 typedef union { __m256 data; } float8;
870 typedef union { __m256 data[2]; } float16;
871 
872 typedef union { double data; } double1;
873 typedef union { __m128d data; } double2;
874 typedef union { __m256d data; } double4;
875 typedef union { double4 data; } double3;
876 typedef union { __m256d data[2]; } double8;
877 typedef union { __m256d data[4]; } double16;
878 
879 #endif // defined(_MSC_VER)
880 #endif
TODO-doc.
#define __host__
Definition: host_defines.h:41