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  inline __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  explicit inline __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  inline __host__ __device__
127  HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast<T>(xs)...}; }
128  inline __host__ __device__
129  HIP_vector_type(const HIP_vector_type&) = default;
130  inline __host__ __device__
131  HIP_vector_type(HIP_vector_type&&) = default;
132  inline __host__ __device__
133  ~HIP_vector_type() = default;
134 
135  inline __host__ __device__
136  HIP_vector_type& operator=(const HIP_vector_type&) = default;
137  inline __host__ __device__
138  HIP_vector_type& operator=(HIP_vector_type&&) = default;
139 
140  // Operators
141  inline __host__ __device__
142  HIP_vector_type& operator++() noexcept
143  {
144  return *this += HIP_vector_type{1};
145  }
146  inline __host__ __device__
147  HIP_vector_type operator++(int) noexcept
148  {
149  auto tmp(*this);
150  ++*this;
151  return tmp;
152  }
153  inline __host__ __device__
154  HIP_vector_type& operator--() noexcept
155  {
156  return *this -= HIP_vector_type{1};
157  }
158  inline __host__ __device__
159  HIP_vector_type operator--(int) noexcept
160  {
161  auto tmp(*this);
162  --*this;
163  return tmp;
164  }
165  inline __host__ __device__
166  HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
167  {
168  data += x.data;
169  return *this;
170  }
171  inline __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  inline __host__ __device__
182  HIP_vector_type& operator-=(U x) noexcept
183  {
184  return *this -= HIP_vector_type{x};
185  }
186  inline __host__ __device__
187  HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
188  {
189  data *= x.data;
190  return *this;
191  }
192  inline __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  inline __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  inline __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  inline __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  inline __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  inline __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  inline __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  inline __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  inline __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  inline __host__ __device__
279  HIP_vector_type<T, n> operator+(
280  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
281  {
282  return HIP_vector_type<T, n>{x} += y;
283  }
284  template<typename T, unsigned int n, typename U>
285  inline __host__ __device__
286  HIP_vector_type<T, n> operator+(
287  const HIP_vector_type<T, n>& x, U y) noexcept
288  {
289  return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
290  }
291  template<typename T, unsigned int n, typename U>
292  inline __host__ __device__
293  HIP_vector_type<T, n> operator+(
294  U x, const HIP_vector_type<T, n>& y) noexcept
295  {
296  return HIP_vector_type<T, n>{x} += y;
297  }
298 
299  template<typename T, unsigned int n>
300  inline __host__ __device__
301  HIP_vector_type<T, n> operator-(
302  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
303  {
304  return HIP_vector_type<T, n>{x} -= y;
305  }
306  template<typename T, unsigned int n, typename U>
307  inline __host__ __device__
308  HIP_vector_type<T, n> operator-(
309  const HIP_vector_type<T, n>& x, U y) noexcept
310  {
311  return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
312  }
313  template<typename T, unsigned int n, typename U>
314  inline __host__ __device__
315  HIP_vector_type<T, n> operator-(
316  U x, const HIP_vector_type<T, n>& y) noexcept
317  {
318  return HIP_vector_type<T, n>{x} -= y;
319  }
320 
321  template<typename T, unsigned int n>
322  inline __host__ __device__
323  HIP_vector_type<T, n> operator*(
324  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
325  {
326  return HIP_vector_type<T, n>{x} *= y;
327  }
328  template<typename T, unsigned int n, typename U>
329  inline __host__ __device__
330  HIP_vector_type<T, n> operator*(
331  const HIP_vector_type<T, n>& x, U y) noexcept
332  {
333  return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
334  }
335  template<typename T, unsigned int n, typename U>
336  inline __host__ __device__
337  HIP_vector_type<T, n> operator*(
338  U x, const HIP_vector_type<T, n>& y) noexcept
339  {
340  return HIP_vector_type<T, n>{x} *= y;
341  }
342 
343  template<typename T, unsigned int n>
344  inline __host__ __device__
345  HIP_vector_type<T, n> operator/(
346  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
347  {
348  return HIP_vector_type<T, n>{x} /= y;
349  }
350  template<typename T, unsigned int n, typename U>
351  inline __host__ __device__
352  HIP_vector_type<T, n> operator/(
353  const HIP_vector_type<T, n>& x, U y) noexcept
354  {
355  return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
356  }
357  template<typename T, unsigned int n, typename U>
358  inline __host__ __device__
359  HIP_vector_type<T, n> operator/(
360  U x, const HIP_vector_type<T, n>& y) noexcept
361  {
362  return HIP_vector_type<T, n>{x} /= y;
363  }
364 
365  template<typename T, unsigned int n>
366  inline __host__ __device__
367  bool operator==(
368  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
369  {
370  auto tmp = x.data == y.data;
371  for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false;
372  return true;
373  }
374  template<typename T, unsigned int n, typename U>
375  inline __host__ __device__
376  bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
377  {
378  return x == HIP_vector_type<T, n>{y};
379  }
380  template<typename T, unsigned int n, typename U>
381  inline __host__ __device__
382  bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
383  {
384  return HIP_vector_type<T, n>{x} == y;
385  }
386 
387  template<typename T, unsigned int n>
388  inline __host__ __device__
389  bool operator!=(
390  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
391  {
392  return !(x == y);
393  }
394  template<typename T, unsigned int n, typename U>
395  inline __host__ __device__
396  bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
397  {
398  return !(x == y);
399  }
400  template<typename T, unsigned int n, typename U>
401  inline __host__ __device__
402  bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
403  {
404  return !(x == y);
405  }
406 
407  template<
408  typename T,
409  unsigned int n,
410  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
411  inline __host__ __device__
412  HIP_vector_type<T, n> operator%(
413  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
414  {
415  return HIP_vector_type<T, n>{x} %= y;
416  }
417  template<
418  typename T,
419  unsigned int n,
420  typename U,
421  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
422  inline __host__ __device__
423  HIP_vector_type<T, n> operator%(
424  const HIP_vector_type<T, n>& x, U y) noexcept
425  {
426  return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
427  }
428  template<
429  typename T,
430  unsigned int n,
431  typename U,
432  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
433  inline __host__ __device__
434  HIP_vector_type<T, n> operator%(
435  U x, const HIP_vector_type<T, n>& y) noexcept
436  {
437  return HIP_vector_type<T, n>{x} %= y;
438  }
439 
440  template<
441  typename T,
442  unsigned int n,
443  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
444  inline __host__ __device__
445  HIP_vector_type<T, n> operator^(
446  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
447  {
448  return HIP_vector_type<T, n>{x} ^= y;
449  }
450  template<
451  typename T,
452  unsigned int n,
453  typename U,
454  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
455  inline __host__ __device__
456  HIP_vector_type<T, n> operator^(
457  const HIP_vector_type<T, n>& x, U y) noexcept
458  {
459  return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
460  }
461  template<
462  typename T,
463  unsigned int n,
464  typename U,
465  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
466  inline __host__ __device__
467  HIP_vector_type<T, n> operator^(
468  U x, const HIP_vector_type<T, n>& y) noexcept
469  {
470  return HIP_vector_type<T, n>{x} ^= y;
471  }
472 
473  template<
474  typename T,
475  unsigned int n,
476  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
477  inline __host__ __device__
478  HIP_vector_type<T, n> operator|(
479  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
480  {
481  return HIP_vector_type<T, n>{x} |= y;
482  }
483  template<
484  typename T,
485  unsigned int n,
486  typename U,
487  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
488  inline __host__ __device__
489  HIP_vector_type<T, n> operator|(
490  const HIP_vector_type<T, n>& x, U y) noexcept
491  {
492  return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
493  }
494  template<
495  typename T,
496  unsigned int n,
497  typename U,
498  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
499  inline __host__ __device__
500  HIP_vector_type<T, n> operator|(
501  U x, const HIP_vector_type<T, n>& y) noexcept
502  {
503  return HIP_vector_type<T, n>{x} |= y;
504  }
505 
506  template<
507  typename T,
508  unsigned int n,
509  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
510  inline __host__ __device__
511  HIP_vector_type<T, n> operator&(
512  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
513  {
514  return HIP_vector_type<T, n>{x} &= y;
515  }
516  template<
517  typename T,
518  unsigned int n,
519  typename U,
520  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
521  inline __host__ __device__
522  HIP_vector_type<T, n> operator&(
523  const HIP_vector_type<T, n>& x, U y) noexcept
524  {
525  return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
526  }
527  template<
528  typename T,
529  unsigned int n,
530  typename U,
531  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
532  inline __host__ __device__
533  HIP_vector_type<T, n> operator&(
534  U x, const HIP_vector_type<T, n>& y) noexcept
535  {
536  return HIP_vector_type<T, n>{x} &= y;
537  }
538 
539  template<
540  typename T,
541  unsigned int n,
542  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
543  inline __host__ __device__
544  HIP_vector_type<T, n> operator>>(
545  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
546  {
547  return HIP_vector_type<T, n>{x} >>= y;
548  }
549  template<
550  typename T,
551  unsigned int n,
552  typename U,
553  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
554  inline __host__ __device__
555  HIP_vector_type<T, n> operator>>(
556  const HIP_vector_type<T, n>& x, U y) noexcept
557  {
558  return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
559  }
560  template<
561  typename T,
562  unsigned int n,
563  typename U,
564  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
565  inline __host__ __device__
566  HIP_vector_type<T, n> operator>>(
567  U x, const HIP_vector_type<T, n>& y) noexcept
568  {
569  return HIP_vector_type<T, n>{x} >>= y;
570  }
571 
572  template<
573  typename T,
574  unsigned int n,
575  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
576  inline __host__ __device__
577  HIP_vector_type<T, n> operator<<(
578  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
579  {
580  return HIP_vector_type<T, n>{x} <<= y;
581  }
582  template<
583  typename T,
584  unsigned int n,
585  typename U,
586  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
587  inline __host__ __device__
588  HIP_vector_type<T, n> operator<<(
589  const HIP_vector_type<T, n>& x, U y) noexcept
590  {
591  return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
592  }
593  template<
594  typename T,
595  unsigned int n,
596  typename U,
597  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
598  inline __host__ __device__
599  HIP_vector_type<T, n> operator<<(
600  U x, const HIP_vector_type<T, n>& y) noexcept
601  {
602  return HIP_vector_type<T, n>{x} <<= y;
603  }
604 
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>;
610 #else
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);\
616  typedef struct {\
617  union {\
618  CUDA_name##_impl1 data;\
619  struct {\
620  T x;\
621  };\
622  };\
623  } CUDA_name##1;\
624  typedef struct {\
625  union {\
626  CUDA_name##_impl2 data;\
627  struct {\
628  T x;\
629  T y;\
630  };\
631  };\
632  } CUDA_name##2;\
633  typedef struct {\
634  union {\
635  CUDA_name##_impl3 data;\
636  struct {\
637  T x;\
638  T y;\
639  T z;\
640  };\
641  };\
642  } CUDA_name##3;\
643  typedef struct {\
644  union {\
645  CUDA_name##_impl4 data;\
646  struct {\
647  T x;\
648  T y;\
649  T z;\
650  T w;\
651  };\
652  };\
653  } CUDA_name##4;
654 #endif
655 
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);
668 
669 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
670  static inline __device__ __host__ \
671  type make_##type(comp x) { type r{x}; return r; }
672 
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; }
676 
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; }
680 
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}; \
685  return r; \
686  }
687 
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);
692 
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);
697 
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);
702 
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);
707 
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);
712 
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);
717 
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);
722 
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);
727 
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);
732 
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);
737 
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);
742 
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>
752 
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;
759 
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;
766 
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;
773 
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;
780 
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;
787 
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;
794 
795 #if !defined(_WIN64)
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;
802 
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;
816 
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)
824 
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;
831 
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;
838 
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;
845 
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;
852 
853 #endif // defined(_MSC_VER)
854 #endif
TODO-doc.
#define __host__
Definition: host_defines.h:41