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 
35 #include "hip/hcc_detail/host_defines.h"
36 
37 #if defined(__clang__)
38  #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
39 #elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax.
40  #define __ROUND_UP_TO_NEXT_POT__(x) \
41  (1 << (31 - __builtin_clz(x) + (x > (1 << (31 - __builtin_clz(x))))))
42  #define __NATIVE_VECTOR__(n, T) \
43  __attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T))))
44 #endif
45 
46 #if defined(__cplusplus)
47  #include <type_traits>
48 
49  template<typename T, unsigned int n> struct HIP_vector_base;
50 
51  template<typename T>
52  struct HIP_vector_base<T, 1> {
53  typedef T Native_vec_ __NATIVE_VECTOR__(1, T);
54 
55  union {
56  Native_vec_ data;
57  struct {
58  T x;
59  };
60  };
61  };
62 
63  template<typename T>
64  struct HIP_vector_base<T, 2> {
65  typedef T Native_vec_ __NATIVE_VECTOR__(2, T);
66 
67  union {
68  Native_vec_ data;
69  struct {
70  T x;
71  T y;
72  };
73  };
74  };
75 
76  template<typename T>
77  struct HIP_vector_base<T, 3> {
78  typedef T Native_vec_ __NATIVE_VECTOR__(3, T);
79 
80  union {
81  Native_vec_ data;
82  struct {
83  T x;
84  T y;
85  T z;
86  };
87  };
88  };
89 
90  template<typename T>
91  struct HIP_vector_base<T, 4> {
92  typedef T Native_vec_ __NATIVE_VECTOR__(4, T);
93 
94  union {
95  Native_vec_ data;
96  struct {
97  T x;
98  T y;
99  T z;
100  T w;
101  };
102  };
103  };
104 
105  template<typename T, unsigned int rank>
106  struct HIP_vector_type : public HIP_vector_base<T, rank> {
107  using HIP_vector_base<T, rank>::data;
108  using typename HIP_vector_base<T, rank>::Native_vec_;
109 
110  __host__ __device__
111  HIP_vector_type() = default;
112  template<
113  typename U,
114  typename std::enable_if<
115  std::is_convertible<U, T>{}>::type* = nullptr>
116  __host__ __device__
117  HIP_vector_type(U x) noexcept
118  {
119  for (auto i = 0u; i != rank; ++i) data[i] = x;
120  }
121  template< // TODO: constrain based on type as well.
122  typename... Us,
123  typename std::enable_if<
124  (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
125  __host__ __device__
126  HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast<T>(xs)...}; }
127  __host__ __device__
128  HIP_vector_type(const HIP_vector_type&) = default;
129  __host__ __device__
130  HIP_vector_type(HIP_vector_type&&) = default;
131  __host__ __device__
132  ~HIP_vector_type() = default;
133 
134  __host__ __device__
135  HIP_vector_type& operator=(const HIP_vector_type&) = default;
136  __host__ __device__
137  HIP_vector_type& operator=(HIP_vector_type&&) = default;
138 
139  // Operators
140  __host__ __device__
141  HIP_vector_type& operator++() noexcept
142  {
143  return *this += HIP_vector_type{1};
144  }
145  __host__ __device__
146  HIP_vector_type operator++(int) noexcept
147  {
148  auto tmp(*this);
149  ++*this;
150  return tmp;
151  }
152  __host__ __device__
153  HIP_vector_type& operator--() noexcept
154  {
155  return *this -= HIP_vector_type{1};
156  }
157  __host__ __device__
158  HIP_vector_type operator--(int) noexcept
159  {
160  auto tmp(*this);
161  --*this;
162  return tmp;
163  }
164  __host__ __device__
165  HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
166  {
167  data += x.data;
168  return *this;
169  }
170  __host__ __device__
171  HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
172  {
173  data -= x.data;
174  return *this;
175  }
176  template<
177  typename U,
178  typename std::enable_if<
179  std::is_convertible<U, T>{}>::type* = nullptr>
180  __host__ __device__
181  HIP_vector_type& operator-=(U x) noexcept
182  {
183  return *this -= HIP_vector_type{x};
184  }
185  __host__ __device__
186  HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
187  {
188  data *= x.data;
189  return *this;
190  }
191  __host__ __device__
192  HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
193  {
194  data /= x.data;
195  return *this;
196  }
197 
198  template<
199  typename U = T,
200  typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
201  __host__ __device__
202  HIP_vector_type operator-() noexcept
203  {
204  auto tmp(*this);
205  tmp.data = -tmp.data;
206  return tmp;
207  }
208 
209  template<
210  typename U = T,
211  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
212  __host__ __device__
213  HIP_vector_type operator~() noexcept
214  {
215  HIP_vector_type r{*this};
216  r.data = ~r.data;
217  return r;
218  }
219  template<
220  typename U = T,
221  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
222  __host__ __device__
223  HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
224  {
225  data %= x.data;
226  return *this;
227  }
228  template<
229  typename U = T,
230  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
231  __host__ __device__
232  HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
233  {
234  data ^= x.data;
235  return *this;
236  }
237  template<
238  typename U = T,
239  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
240  __host__ __device__
241  HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
242  {
243  data |= x.data;
244  return *this;
245  }
246  template<
247  typename U = T,
248  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
249  __host__ __device__
250  HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
251  {
252  data &= x.data;
253  return *this;
254  }
255  template<
256  typename U = T,
257  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
258  __host__ __device__
259  HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
260  {
261  data >>= x.data;
262  return *this;
263  }
264  template<
265  typename U = T,
266  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
267  __host__ __device__
268  HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
269  {
270  data <<= x.data;
271  return *this;
272  }
273  };
274 
275 
276  template<typename T, unsigned int n>
277  __host__ __device__
278  inline
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  __host__ __device__
286  inline
287  HIP_vector_type<T, n> operator+(
288  const HIP_vector_type<T, n>& x, U y) noexcept
289  {
290  return HIP_vector_type<T, n>{x} += y;
291  }
292  template<typename T, unsigned int n, typename U>
293  __host__ __device__
294  inline
295  HIP_vector_type<T, n> operator+(
296  U x, const HIP_vector_type<T, n>& y) noexcept
297  {
298  return y + x;
299  }
300 
301  template<typename T, unsigned int n>
302  __host__ __device__
303  inline
304  HIP_vector_type<T, n> operator-(
305  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
306  {
307  return HIP_vector_type<T, n>{x} -= y;
308  }
309  template<typename T, unsigned int n, typename U>
310  __host__ __device__
311  inline
312  HIP_vector_type<T, n> operator-(
313  const HIP_vector_type<T, n>& x, U y) noexcept
314  {
315  return HIP_vector_type<T, n>{x} -= y;
316  }
317  template<typename T, unsigned int n, typename U>
318  __host__ __device__
319  inline
320  HIP_vector_type<T, n> operator-(
321  U x, const HIP_vector_type<T, n>& y) noexcept
322  {
323  return HIP_vector_type<T, n>{x} -= y;
324  }
325 
326  template<typename T, unsigned int n>
327  __host__ __device__
328  inline
329  HIP_vector_type<T, n> operator*(
330  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
331  {
332  return HIP_vector_type<T, n>{x} *= y;
333  }
334  template<typename T, unsigned int n, typename U>
335  __host__ __device__
336  inline
337  HIP_vector_type<T, n> operator*(
338  const HIP_vector_type<T, n>& x, U y) noexcept
339  {
340  return HIP_vector_type<T, n>{x} *= y;
341  }
342  template<typename T, unsigned int n, typename U>
343  __host__ __device__
344  inline
345  HIP_vector_type<T, n> operator*(
346  U x, const HIP_vector_type<T, n>& y) noexcept
347  {
348  return y * x;
349  }
350 
351  template<typename T, unsigned int n>
352  __host__ __device__
353  inline
354  HIP_vector_type<T, n> operator/(
355  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
356  {
357  return HIP_vector_type<T, n>{x} /= y;
358  }
359  template<typename T, unsigned int n, typename U>
360  __host__ __device__
361  inline
362  HIP_vector_type<T, n> operator/(
363  const HIP_vector_type<T, n>& x, U y) noexcept
364  {
365  return HIP_vector_type<T, n>{x} /= y;
366  }
367  template<typename T, unsigned int n, typename U>
368  __host__ __device__
369  inline
370  HIP_vector_type<T, n> operator/(
371  U x, const HIP_vector_type<T, n>& y) noexcept
372  {
373  return HIP_vector_type<T, n>{x} /= y;
374  }
375 
376  template<typename T, unsigned int n>
377  __host__ __device__
378  inline
379  bool operator==(
380  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
381  {
382  auto tmp = x.data == y.data;
383  for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false;
384  return true;
385  }
386  template<typename T, unsigned int n, typename U>
387  __host__ __device__
388  inline
389  bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
390  {
391  return x == HIP_vector_type<T, n>{y};
392  }
393  template<typename T, unsigned int n, typename U>
394  __host__ __device__
395  inline
396  bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
397  {
398  return HIP_vector_type<T, n>{x} == y;
399  }
400 
401  template<typename T, unsigned int n>
402  __host__ __device__
403  inline
404  bool operator!=(
405  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
406  {
407  return !(x == y);
408  }
409  template<typename T, unsigned int n, typename U>
410  __host__ __device__
411  inline
412  bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
413  {
414  return !(x == y);
415  }
416  template<typename T, unsigned int n, typename U>
417  __host__ __device__
418  inline
419  bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
420  {
421  return !(x == y);
422  }
423 
424  template<
425  typename T,
426  unsigned int n,
427  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
428  inline
429  HIP_vector_type<T, n> operator%(
430  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
431  {
432  return HIP_vector_type<T, n>{x} %= y;
433  }
434  template<
435  typename T,
436  unsigned int n,
437  typename U,
438  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
439  inline
440  HIP_vector_type<T, n> operator%(
441  const HIP_vector_type<T, n>& x, U y) noexcept
442  {
443  return HIP_vector_type<T, n>{x} %= y;
444  }
445  template<
446  typename T,
447  unsigned int n,
448  typename U,
449  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
450  inline
451  HIP_vector_type<T, n> operator%(
452  U x, const HIP_vector_type<T, n>& y) noexcept
453  {
454  return HIP_vector_type<T, n>{x} %= y;
455  }
456 
457  template<
458  typename T,
459  unsigned int n,
460  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
461  inline
462  HIP_vector_type<T, n> operator^(
463  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
464  {
465  return HIP_vector_type<T, n>{x} ^= y;
466  }
467  template<
468  typename T,
469  unsigned int n,
470  typename U,
471  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
472  inline
473  HIP_vector_type<T, n> operator^(
474  const HIP_vector_type<T, n>& x, U y) noexcept
475  {
476  return HIP_vector_type<T, n>{x} ^= y;
477  }
478  template<
479  typename T,
480  unsigned int n,
481  typename U,
482  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
483  inline
484  HIP_vector_type<T, n> operator^(
485  U x, const HIP_vector_type<T, n>& y) noexcept
486  {
487  return HIP_vector_type<T, n>{x} ^= y;
488  }
489 
490  template<
491  typename T,
492  unsigned int n,
493  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
494  inline
495  HIP_vector_type<T, n> operator|(
496  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
497  {
498  return HIP_vector_type<T, n>{x} |= y;
499  }
500  template<
501  typename T,
502  unsigned int n,
503  typename U,
504  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
505  inline
506  HIP_vector_type<T, n> operator|(
507  const HIP_vector_type<T, n>& x, U y) noexcept
508  {
509  return HIP_vector_type<T, n>{x} |= y;
510  }
511  template<
512  typename T,
513  unsigned int n,
514  typename U,
515  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
516  inline
517  HIP_vector_type<T, n> operator|(
518  U x, const HIP_vector_type<T, n>& y) noexcept
519  {
520  return HIP_vector_type<T, n>{x} |= y;
521  }
522 
523  template<
524  typename T,
525  unsigned int n,
526  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
527  inline
528  HIP_vector_type<T, n> operator&(
529  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
530  {
531  return HIP_vector_type<T, n>{x} &= y;
532  }
533  template<
534  typename T,
535  unsigned int n,
536  typename U,
537  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
538  inline
539  HIP_vector_type<T, n> operator&(
540  const HIP_vector_type<T, n>& x, U y) noexcept
541  {
542  return HIP_vector_type<T, n>{x} &= y;
543  }
544  template<
545  typename T,
546  unsigned int n,
547  typename U,
548  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
549  inline
550  HIP_vector_type<T, n> operator&(
551  U x, const HIP_vector_type<T, n>& y) noexcept
552  {
553  return HIP_vector_type<T, n>{x} &= y;
554  }
555 
556  template<
557  typename T,
558  unsigned int n,
559  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
560  inline
561  HIP_vector_type<T, n> operator>>(
562  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
563  {
564  return HIP_vector_type<T, n>{x} >>= y;
565  }
566  template<
567  typename T,
568  unsigned int n,
569  typename U,
570  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
571  inline
572  HIP_vector_type<T, n> operator>>(
573  const HIP_vector_type<T, n>& x, U y) noexcept
574  {
575  return HIP_vector_type<T, n>{x} >>= y;
576  }
577  template<
578  typename T,
579  unsigned int n,
580  typename U,
581  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
582  inline
583  HIP_vector_type<T, n> operator>>(
584  U x, const HIP_vector_type<T, n>& y) noexcept
585  {
586  return HIP_vector_type<T, n>{x} >>= y;
587  }
588 
589  template<
590  typename T,
591  unsigned int n,
592  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
593  inline
594  HIP_vector_type<T, n> operator<<(
595  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
596  {
597  return HIP_vector_type<T, n>{x} <<= y;
598  }
599  template<
600  typename T,
601  unsigned int n,
602  typename U,
603  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
604  inline
605  HIP_vector_type<T, n> operator<<(
606  const HIP_vector_type<T, n>& x, U y) noexcept
607  {
608  return HIP_vector_type<T, n>{x} <<= y;
609  }
610  template<
611  typename T,
612  unsigned int n,
613  typename U,
614  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
615  inline
616  HIP_vector_type<T, n> operator<<(
617  U x, const HIP_vector_type<T, n>& y) noexcept
618  {
619  return HIP_vector_type<T, n>{x} <<= y;
620  }
621 
622  #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
623  using CUDA_name##1 = HIP_vector_type<T, 1>;\
624  using CUDA_name##2 = HIP_vector_type<T, 2>;\
625  using CUDA_name##3 = HIP_vector_type<T, 3>;\
626  using CUDA_name##4 = HIP_vector_type<T, 4>;
627 #else
628  #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
629  typedef T CUDA_name##_impl1 __NATIVE_VECTOR__(1, T);\
630  typedef T CUDA_name##_impl2 __NATIVE_VECTOR__(2, T);\
631  typedef T CUDA_name##_impl3 __NATIVE_VECTOR__(3, T);\
632  typedef T CUDA_name##_impl4 __NATIVE_VECTOR__(4, T);\
633  typedef struct {\
634  union {\
635  CUDA_name##_impl1 data;\
636  struct {\
637  T x;\
638  };\
639  };\
640  } CUDA_name##1;\
641  typedef struct {\
642  union {\
643  CUDA_name##_impl2 data;\
644  struct {\
645  T x;\
646  T y;\
647  };\
648  };\
649  } CUDA_name##2;\
650  typedef struct {\
651  union {\
652  CUDA_name##_impl3 data;\
653  struct {\
654  T x;\
655  T y;\
656  T z;\
657  };\
658  };\
659  } CUDA_name##3;\
660  typedef struct {\
661  union {\
662  CUDA_name##_impl4 data;\
663  struct {\
664  T x;\
665  T y;\
666  T z;\
667  T w;\
668  };\
669  };\
670  } CUDA_name##4;
671 #endif
672 
673 __MAKE_VECTOR_TYPE__(uchar, unsigned char);
674 __MAKE_VECTOR_TYPE__(char, char);
675 __MAKE_VECTOR_TYPE__(ushort, unsigned short);
676 __MAKE_VECTOR_TYPE__(short, short);
677 __MAKE_VECTOR_TYPE__(uint, unsigned int);
678 __MAKE_VECTOR_TYPE__(int, int);
679 __MAKE_VECTOR_TYPE__(ulong, unsigned long);
680 __MAKE_VECTOR_TYPE__(long, long);
681 __MAKE_VECTOR_TYPE__(ulonglong, unsigned long long);
682 __MAKE_VECTOR_TYPE__(longlong, long long);
683 __MAKE_VECTOR_TYPE__(float, float);
684 __MAKE_VECTOR_TYPE__(double, double);
685 
686 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
687  __device__ __host__ \
688  static \
689  inline \
690  type make_##type(comp x) { type r = {x}; return r; }
691 
692 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
693  __device__ __host__ \
694  static \
695  inline \
696  type make_##type(comp x, comp y) { type r = {x, y}; return r; }
697 
698 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
699  __device__ __host__ \
700  static \
701  inline \
702  type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; }
703 
704 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
705  __device__ __host__ \
706  static \
707  inline \
708  type make_##type(comp x, comp y, comp z, comp w) { \
709  type r = {x, y, z, w}; \
710  return r; \
711  }
712 
713 DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1);
714 DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2);
715 DECLOP_MAKE_THREE_COMPONENT(unsigned char, uchar3);
716 DECLOP_MAKE_FOUR_COMPONENT(unsigned char, uchar4);
717 
718 DECLOP_MAKE_ONE_COMPONENT(signed char, char1);
719 DECLOP_MAKE_TWO_COMPONENT(signed char, char2);
720 DECLOP_MAKE_THREE_COMPONENT(signed char, char3);
721 DECLOP_MAKE_FOUR_COMPONENT(signed char, char4);
722 
723 DECLOP_MAKE_ONE_COMPONENT(unsigned short, ushort1);
724 DECLOP_MAKE_TWO_COMPONENT(unsigned short, ushort2);
725 DECLOP_MAKE_THREE_COMPONENT(unsigned short, ushort3);
726 DECLOP_MAKE_FOUR_COMPONENT(unsigned short, ushort4);
727 
728 DECLOP_MAKE_ONE_COMPONENT(signed short, short1);
729 DECLOP_MAKE_TWO_COMPONENT(signed short, short2);
730 DECLOP_MAKE_THREE_COMPONENT(signed short, short3);
731 DECLOP_MAKE_FOUR_COMPONENT(signed short, short4);
732 
733 DECLOP_MAKE_ONE_COMPONENT(unsigned int, uint1);
734 DECLOP_MAKE_TWO_COMPONENT(unsigned int, uint2);
735 DECLOP_MAKE_THREE_COMPONENT(unsigned int, uint3);
736 DECLOP_MAKE_FOUR_COMPONENT(unsigned int, uint4);
737 
738 DECLOP_MAKE_ONE_COMPONENT(signed int, int1);
739 DECLOP_MAKE_TWO_COMPONENT(signed int, int2);
740 DECLOP_MAKE_THREE_COMPONENT(signed int, int3);
741 DECLOP_MAKE_FOUR_COMPONENT(signed int, int4);
742 
743 DECLOP_MAKE_ONE_COMPONENT(float, float1);
744 DECLOP_MAKE_TWO_COMPONENT(float, float2);
745 DECLOP_MAKE_THREE_COMPONENT(float, float3);
746 DECLOP_MAKE_FOUR_COMPONENT(float, float4);
747 
748 DECLOP_MAKE_ONE_COMPONENT(double, double1);
749 DECLOP_MAKE_TWO_COMPONENT(double, double2);
750 DECLOP_MAKE_THREE_COMPONENT(double, double3);
751 DECLOP_MAKE_FOUR_COMPONENT(double, double4);
752 
753 DECLOP_MAKE_ONE_COMPONENT(unsigned long, ulong1);
754 DECLOP_MAKE_TWO_COMPONENT(unsigned long, ulong2);
755 DECLOP_MAKE_THREE_COMPONENT(unsigned long, ulong3);
756 DECLOP_MAKE_FOUR_COMPONENT(unsigned long, ulong4);
757 
758 DECLOP_MAKE_ONE_COMPONENT(signed long, long1);
759 DECLOP_MAKE_TWO_COMPONENT(signed long, long2);
760 DECLOP_MAKE_THREE_COMPONENT(signed long, long3);
761 DECLOP_MAKE_FOUR_COMPONENT(signed long, long4);
762 
763 DECLOP_MAKE_ONE_COMPONENT(unsigned long long, ulonglong1);
764 DECLOP_MAKE_TWO_COMPONENT(unsigned long long, ulonglong2);
765 DECLOP_MAKE_THREE_COMPONENT(unsigned long long, ulonglong3);
766 DECLOP_MAKE_FOUR_COMPONENT(unsigned long long, ulonglong4);
767 
768 DECLOP_MAKE_ONE_COMPONENT(signed long long, longlong1);
769 DECLOP_MAKE_TWO_COMPONENT(signed long long, longlong2);
770 DECLOP_MAKE_THREE_COMPONENT(signed long long, longlong3);
771 DECLOP_MAKE_FOUR_COMPONENT(signed long long, longlong4);
772 
773 #endif
host_defines.h
TODO-doc.
__host__
#define __host__
Definition: host_defines.h:41
Generated on Thu Dec 13 2018 02:17:13 for HIP: Heterogenous-computing Interface for Portability by   doxygen 1.8.11