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_AMD_DETAIL_HIP_VECTOR_TYPES_H
29 #define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_VECTOR_TYPES_H
30 
32 
33 #if defined(__HIPCC_RTC__)
34  #define __HOST_DEVICE__ __device__
35 #else
36  #define __HOST_DEVICE__ __host__ __device__
37 #endif
38 
39 #if defined(__has_attribute)
40  #if __has_attribute(ext_vector_type)
41  #define __NATIVE_VECTOR__(n, T) T __attribute__((ext_vector_type(n)))
42  #else
43  #define __NATIVE_VECTOR__(n, T) T[n]
44  #endif
45 
46 #if defined(__cplusplus)
47 #if !defined(__HIPCC_RTC__)
48  #include <array>
49  #include <iosfwd>
50  #include <type_traits>
51 #endif // !defined(__HIPCC_RTC__)
52 
53  namespace hip_impl {
54  template<typename, typename, unsigned int> struct Scalar_accessor;
55  } // Namespace hip_impl.
56 
57  namespace std {
58  template<typename T, typename U, unsigned int n>
59  struct is_integral<hip_impl::Scalar_accessor<T, U, n>>
60  : is_integral<T> {};
61  template<typename T, typename U, unsigned int n>
62  struct is_floating_point<hip_impl::Scalar_accessor<T, U, n>>
63  : is_floating_point<T> {};
64  } // Namespace std.
65 
66  namespace hip_impl {
67  template<typename T, typename Vector, unsigned int idx>
68  struct Scalar_accessor {
69  struct Address {
70  const Scalar_accessor* p;
71 
72  __HOST_DEVICE__
73  operator const T*() const noexcept {
74  return &reinterpret_cast<const T*>(p)[idx];
75  }
76  __HOST_DEVICE__
77  operator const T*() const volatile noexcept {
78  return &reinterpret_cast<const T*>(p)[idx];
79  }
80  __HOST_DEVICE__
81  operator T*() noexcept {
82  return &reinterpret_cast<T*>(
83  const_cast<Scalar_accessor*>(p))[idx];
84  }
85  __HOST_DEVICE__
86  operator T*() volatile noexcept {
87  return &reinterpret_cast<T*>(
88  const_cast<Scalar_accessor*>(p))[idx];
89  }
90  };
91 
92  friend
93  inline
94  std::ostream& operator<<(std::ostream& os,
95  const Scalar_accessor& x) noexcept {
96  return os << x.data[idx];
97  }
98  friend
99  inline
100  std::istream& operator>>(std::istream& is,
101  Scalar_accessor& x) noexcept {
102  T tmp;
103  is >> tmp;
104  x.data[idx] = tmp;
105 
106  return is;
107  }
108 
109  // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor
110  Vector data;
111 
112  __HOST_DEVICE__
113  operator T() const noexcept { return data[idx]; }
114  __HOST_DEVICE__
115  operator T() const volatile noexcept { return data[idx]; }
116 
117 #ifdef __HIP_ENABLE_VECTOR_SCALAR_ACCESSORY_ENUM_CONVERSION__
118  // The conversions to enum are fairly ghastly, but unfortunately used in
119  // some pre-existing, difficult to modify, code.
120  template<
121  typename U,
122  typename std::enable_if<
123  !std::is_same<U, T>{} &&
124  std::is_enum<U>{} &&
125  std::is_convertible<
126  T, typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* = nullptr>
127  __HOST_DEVICE__
128  operator U() const noexcept { return static_cast<U>(data[idx]); }
129  template<
130  typename U,
131  typename std::enable_if<
132  !std::is_same<U, T>{} &&
133  std::is_enum<U>{} &&
134  std::is_convertible<
135  T, typename std::enable_if<std::is_enum<U>::value, std::underlying_type<U>>::type::type>{}>::type* = nullptr>
136  __HOST_DEVICE__
137  operator U() const volatile noexcept { return static_cast<U>(data[idx]); }
138 #endif
139 
140  __HOST_DEVICE__
141  operator T&() noexcept {
142  return reinterpret_cast<
143  T (&)[sizeof(Vector) / sizeof(T)]>(data)[idx];
144  }
145  __HOST_DEVICE__
146  operator volatile T&() volatile noexcept {
147  return reinterpret_cast<
148  volatile T (&)[sizeof(Vector) / sizeof(T)]>(data)[idx];
149  }
150 
151  __HOST_DEVICE__
152  Address operator&() const noexcept { return Address{this}; }
153 
154  __HOST_DEVICE__
155  Scalar_accessor& operator=(const Scalar_accessor& x) noexcept {
156  data[idx] = x.data[idx];
157 
158  return *this;
159  }
160  __HOST_DEVICE__
161  Scalar_accessor& operator=(T x) noexcept {
162  data[idx] = x;
163 
164  return *this;
165  }
166  __HOST_DEVICE__
167  volatile Scalar_accessor& operator=(T x) volatile noexcept {
168  data[idx] = x;
169 
170  return *this;
171  }
172 
173  __HOST_DEVICE__
174  Scalar_accessor& operator++() noexcept {
175  ++data[idx];
176  return *this;
177  }
178  __HOST_DEVICE__
179  T operator++(int) noexcept {
180  auto r{data[idx]};
181  ++data[idx];
182  return *this;
183  }
184  __HOST_DEVICE__
185  Scalar_accessor& operator--() noexcept {
186  --data[idx];
187  return *this;
188  }
189  __HOST_DEVICE__
190  T operator--(int) noexcept {
191  auto r{data[idx]};
192  --data[idx];
193  return *this;
194  }
195 
196  // TODO: convertibility is too restrictive, constraint should be on
197  // the operator being invocable with a value of type U.
198  template<
199  typename U,
200  typename std::enable_if<
201  std::is_convertible<U, T>{}>::type* = nullptr>
202  __HOST_DEVICE__
203  Scalar_accessor& operator+=(U x) noexcept {
204  data[idx] += x;
205  return *this;
206  }
207  template<
208  typename U,
209  typename std::enable_if<
210  std::is_convertible<U, T>{}>::type* = nullptr>
211  __HOST_DEVICE__
212  Scalar_accessor& operator-=(U x) noexcept {
213  data[idx] -= x;
214  return *this;
215  }
216 
217  template<
218  typename U,
219  typename std::enable_if<
220  std::is_convertible<U, T>{}>::type* = nullptr>
221  __HOST_DEVICE__
222  Scalar_accessor& operator*=(U x) noexcept {
223  data[idx] *= x;
224  return *this;
225  }
226  template<
227  typename U,
228  typename std::enable_if<
229  std::is_convertible<U, T>{}>::type* = nullptr>
230  __HOST_DEVICE__
231  Scalar_accessor& operator/=(U x) noexcept {
232  data[idx] /= x;
233  return *this;
234  }
235  template<
236  typename U = T,
237  typename std::enable_if<std::is_convertible<U, T>{} &&
238  std::is_integral<U>{}>::type* = nullptr>
239  __HOST_DEVICE__
240  Scalar_accessor& operator%=(U x) noexcept {
241  data[idx] %= x;
242  return *this;
243  }
244 
245  template<
246  typename U = T,
247  typename std::enable_if<std::is_convertible<U, T>{} &&
248  std::is_integral<U>{}>::type* = nullptr>
249  __HOST_DEVICE__
250  Scalar_accessor& operator>>=(U x) noexcept {
251  data[idx] >>= x;
252  return *this;
253  }
254  template<
255  typename U = T,
256  typename std::enable_if<std::is_convertible<U, T>{} &&
257  std::is_integral<U>{}>::type* = nullptr>
258  __HOST_DEVICE__
259  Scalar_accessor& operator<<=(U x) noexcept {
260  data[idx] <<= x;
261  return *this;
262  }
263  template<
264  typename U = T,
265  typename std::enable_if<std::is_convertible<U, T>{} &&
266  std::is_integral<U>{}>::type* = nullptr>
267  __HOST_DEVICE__
268  Scalar_accessor& operator&=(U x) noexcept {
269  data[idx] &= x;
270  return *this;
271  }
272  template<
273  typename U = T,
274  typename std::enable_if<std::is_convertible<U, T>{} &&
275  std::is_integral<U>{}>::type* = nullptr>
276  __HOST_DEVICE__
277  Scalar_accessor& operator|=(U x) noexcept {
278  data[idx] |= x;
279  return *this;
280  }
281  template<
282  typename U = T,
283  typename std::enable_if<std::is_convertible<U, T>{} &&
284  std::is_integral<U>{}>::type* = nullptr>
285  __HOST_DEVICE__
286  Scalar_accessor& operator^=(U x) noexcept {
287  data[idx] ^= x;
288  return *this;
289  }
290  };
291 
292  inline
293  constexpr
294  unsigned int next_pot(unsigned int x) {
295  // Precondition: x > 1.
296  return 1u << (32u - __builtin_clz(x - 1u));
297  }
298  } // Namespace hip_impl.
299 
300  template<typename T, unsigned int n> struct HIP_vector_base;
301 
302  template<typename T>
303  struct HIP_vector_base<T, 1> {
304  using Native_vec_ = __NATIVE_VECTOR__(1, T);
305 
306  union {
307  Native_vec_ data;
308 #if __HIP_CLANG_ONLY__
309  struct {
310  T x;
311  };
312 #else
313  hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
314 #endif
315  };
316 
317  using value_type = T;
318 
319  __HOST_DEVICE__
320  HIP_vector_base() = default;
321  __HOST_DEVICE__
322  explicit
323  constexpr
324  HIP_vector_base(T x_) noexcept : data{x_} {}
325  __HOST_DEVICE__
326  constexpr
327  HIP_vector_base(const HIP_vector_base&) = default;
328  __HOST_DEVICE__
329  constexpr
330  HIP_vector_base(HIP_vector_base&&) = default;
331  __HOST_DEVICE__
332  ~HIP_vector_base() = default;
333 
334  __HOST_DEVICE__
335  HIP_vector_base& operator=(const HIP_vector_base& x_) noexcept {
336  #if __has_attribute(ext_vector_type)
337  data = x_.data;
338  #else
339  data[0] = x_.data[0];
340  #endif
341 
342  return *this;
343  }
344  };
345 
346  template<typename T>
347  struct HIP_vector_base<T, 2> {
348  using Native_vec_ = __NATIVE_VECTOR__(2, T);
349 
350  union
351  #if !__has_attribute(ext_vector_type)
352  alignas(hip_impl::next_pot(2 * sizeof(T)))
353  #endif
354  {
355  Native_vec_ data;
356 #if __HIP_CLANG_ONLY__
357  struct {
358  T x;
359  T y;
360  };
361 #else
362  hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
363  hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
364 #endif
365  };
366 
367  using value_type = T;
368 
369  __HOST_DEVICE__
370  HIP_vector_base() = default;
371  __HOST_DEVICE__
372  explicit
373  constexpr
374  HIP_vector_base(T x_) noexcept : data{x_, x_} {}
375  __HOST_DEVICE__
376  constexpr
377  HIP_vector_base(T x_, T y_) noexcept : data{x_, y_} {}
378  __HOST_DEVICE__
379  constexpr
380  HIP_vector_base(const HIP_vector_base&) = default;
381  __HOST_DEVICE__
382  constexpr
383  HIP_vector_base(HIP_vector_base&&) = default;
384  __HOST_DEVICE__
385  ~HIP_vector_base() = default;
386 
387  __HOST_DEVICE__
388  HIP_vector_base& operator=(const HIP_vector_base& x_) noexcept {
389  #if __has_attribute(ext_vector_type)
390  data = x_.data;
391  #else
392  data[0] = x_.data[0];
393  data[1] = x_.data[1];
394  #endif
395 
396  return *this;
397  }
398  };
399 
400  template<typename T>
401  struct HIP_vector_base<T, 3> {
402  struct Native_vec_ {
403  T d[3];
404 
405  __HOST_DEVICE__
406  Native_vec_() = default;
407 
408  __HOST_DEVICE__
409  explicit
410  constexpr
411  Native_vec_(T x_) noexcept : d{x_, x_, x_} {}
412  __HOST_DEVICE__
413  constexpr
414  Native_vec_(T x_, T y_, T z_) noexcept : d{x_, y_, z_} {}
415  __HOST_DEVICE__
416  constexpr
417  Native_vec_(const Native_vec_&) = default;
418  __HOST_DEVICE__
419  constexpr
420  Native_vec_(Native_vec_&&) = default;
421  __HOST_DEVICE__
422  ~Native_vec_() = default;
423 
424  __HOST_DEVICE__
425  Native_vec_& operator=(const Native_vec_&) = default;
426  __HOST_DEVICE__
427  Native_vec_& operator=(Native_vec_&&) = default;
428 
429  __HOST_DEVICE__
430  T& operator[](unsigned int idx) noexcept { return d[idx]; }
431  __HOST_DEVICE__
432  T operator[](unsigned int idx) const noexcept { return d[idx]; }
433 
434  __HOST_DEVICE__
435  Native_vec_& operator+=(const Native_vec_& x_) noexcept
436  {
437  for (auto i = 0u; i != 3u; ++i) d[i] += x_.d[i];
438  return *this;
439  }
440  __HOST_DEVICE__
441  Native_vec_& operator-=(const Native_vec_& x_) noexcept
442  {
443  for (auto i = 0u; i != 3u; ++i) d[i] -= x_.d[i];
444  return *this;
445  }
446 
447  __HOST_DEVICE__
448  Native_vec_& operator*=(const Native_vec_& x_) noexcept
449  {
450  for (auto i = 0u; i != 3u; ++i) d[i] *= x_.d[i];
451  return *this;
452  }
453  __HOST_DEVICE__
454  Native_vec_& operator/=(const Native_vec_& x_) noexcept
455  {
456  for (auto i = 0u; i != 3u; ++i) d[i] /= x_.d[i];
457  return *this;
458  }
459 
460  template<
461  typename U = T,
462  typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
463  __HOST_DEVICE__
464  Native_vec_ operator-() const noexcept
465  {
466  auto r{*this};
467  for (auto&& x : r.d) x = -x;
468  return r;
469  }
470 
471  template<
472  typename U = T,
473  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
474  __HOST_DEVICE__
475  Native_vec_ operator~() const noexcept
476  {
477  auto r{*this};
478  for (auto&& x : r.d) x = ~x;
479  return r;
480  }
481  template<
482  typename U = T,
483  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
484  __HOST_DEVICE__
485  Native_vec_& operator%=(const Native_vec_& x_) noexcept
486  {
487  for (auto i = 0u; i != 3u; ++i) d[i] %= x_.d[i];
488  return *this;
489  }
490  template<
491  typename U = T,
492  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
493  __HOST_DEVICE__
494  Native_vec_& operator^=(const Native_vec_& x_) noexcept
495  {
496  for (auto i = 0u; i != 3u; ++i) d[i] ^= x_.d[i];
497  return *this;
498  }
499  template<
500  typename U = T,
501  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
502  __HOST_DEVICE__
503  Native_vec_& operator|=(const Native_vec_& x_) noexcept
504  {
505  for (auto i = 0u; i != 3u; ++i) d[i] |= x_.d[i];
506  return *this;
507  }
508  template<
509  typename U = T,
510  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
511  __HOST_DEVICE__
512  Native_vec_& operator&=(const Native_vec_& x_) noexcept
513  {
514  for (auto i = 0u; i != 3u; ++i) d[i] &= x_.d[i];
515  return *this;
516  }
517  template<
518  typename U = T,
519  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
520  __HOST_DEVICE__
521  Native_vec_& operator>>=(const Native_vec_& x_) noexcept
522  {
523  for (auto i = 0u; i != 3u; ++i) d[i] >>= x_.d[i];
524  return *this;
525  }
526  template<
527  typename U = T,
528  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
529  __HOST_DEVICE__
530  Native_vec_& operator<<=(const Native_vec_& x_) noexcept
531  {
532  for (auto i = 0u; i != 3u; ++i) d[i] <<= x_.d[i];
533  return *this;
534  }
535 
536  using Vec3_cmp = int __attribute__((vector_size(4 * sizeof(int))));
537  __HOST_DEVICE__
538  Vec3_cmp operator==(const Native_vec_& x_) const noexcept
539  {
540  return Vec3_cmp{d[0] == x_.d[0], d[1] == x_.d[1], d[2] == x_.d[2]};
541  }
542  };
543 
544  union {
545  Native_vec_ data;
546  struct {
547  T x;
548  T y;
549  T z;
550  };
551  };
552 
553  using value_type = T;
554 
555  __HOST_DEVICE__
556  HIP_vector_base() = default;
557  __HOST_DEVICE__
558  explicit
559  constexpr
560  HIP_vector_base(T x_) noexcept : data{x_, x_, x_} {}
561  __HOST_DEVICE__
562  constexpr
563  HIP_vector_base(T x_, T y_, T z_) noexcept : data{x_, y_, z_} {}
564  __HOST_DEVICE__
565  constexpr
566  HIP_vector_base(const HIP_vector_base&) = default;
567  __HOST_DEVICE__
568  constexpr
569  HIP_vector_base(HIP_vector_base&&) = default;
570  __HOST_DEVICE__
571  ~HIP_vector_base() = default;
572 
573  __HOST_DEVICE__
574  HIP_vector_base& operator=(const HIP_vector_base&) = default;
575  __HOST_DEVICE__
576  HIP_vector_base& operator=(HIP_vector_base&&) = default;
577  };
578 
579  template<typename T>
580  struct HIP_vector_base<T, 4> {
581  using Native_vec_ = __NATIVE_VECTOR__(4, T);
582 
583  union
584  #if !__has_attribute(ext_vector_type)
585  alignas(hip_impl::next_pot(4 * sizeof(T)))
586  #endif
587  {
588  Native_vec_ data;
589 #if __HIP_CLANG_ONLY__
590  struct {
591  T x;
592  T y;
593  T z;
594  T w;
595  };
596 #else
597  hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
598  hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
599  hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
600  hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
601 #endif
602  };
603 
604  using value_type = T;
605 
606  __HOST_DEVICE__
607  HIP_vector_base() = default;
608  __HOST_DEVICE__
609  explicit
610  constexpr
611  HIP_vector_base(T x_) noexcept : data{x_, x_, x_, x_} {}
612  __HOST_DEVICE__
613  constexpr
614  HIP_vector_base(T x_, T y_, T z_, T w_) noexcept : data{x_, y_, z_, w_} {}
615  __HOST_DEVICE__
616  constexpr
617  HIP_vector_base(const HIP_vector_base&) = default;
618  __HOST_DEVICE__
619  constexpr
620  HIP_vector_base(HIP_vector_base&&) = default;
621  __HOST_DEVICE__
622  ~HIP_vector_base() = default;
623 
624  __HOST_DEVICE__
625  HIP_vector_base& operator=(const HIP_vector_base& x_) noexcept {
626  #if __has_attribute(ext_vector_type)
627  data = x_.data;
628  #else
629  data[0] = x_.data[0];
630  data[1] = x_.data[1];
631  data[2] = x_.data[2];
632  data[3] = x_.data[3];
633  #endif
634 
635  return *this;
636  }
637  };
638 
639  template<typename T, unsigned int rank>
640  struct HIP_vector_type : public HIP_vector_base<T, rank> {
641  using HIP_vector_base<T, rank>::data;
642  using typename HIP_vector_base<T, rank>::Native_vec_;
643 
644  __HOST_DEVICE__
645  HIP_vector_type() = default;
646  template<
647  typename U,
648  typename std::enable_if<
649  std::is_convertible<U, T>{}>::type* = nullptr>
650  __HOST_DEVICE__
651  explicit
652  constexpr
653  HIP_vector_type(U x_) noexcept
654  : HIP_vector_base<T, rank>{static_cast<T>(x_)}
655  {}
656  template< // TODO: constrain based on type as well.
657  typename... Us,
658  typename std::enable_if<
659  (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
660  __HOST_DEVICE__
661  constexpr
662  HIP_vector_type(Us... xs) noexcept
663  : HIP_vector_base<T, rank>{static_cast<T>(xs)...}
664  {}
665  __HOST_DEVICE__
666  constexpr
667  HIP_vector_type(const HIP_vector_type&) = default;
668  __HOST_DEVICE__
669  constexpr
670  HIP_vector_type(HIP_vector_type&&) = default;
671  __HOST_DEVICE__
672  ~HIP_vector_type() = default;
673 
674  __HOST_DEVICE__
675  HIP_vector_type& operator=(const HIP_vector_type&) = default;
676  __HOST_DEVICE__
677  HIP_vector_type& operator=(HIP_vector_type&&) = default;
678 
679  // Operators
680  __HOST_DEVICE__
681  HIP_vector_type& operator++() noexcept
682  {
683  return *this += HIP_vector_type{1};
684  }
685  __HOST_DEVICE__
686  HIP_vector_type operator++(int) noexcept
687  {
688  auto tmp(*this);
689  ++*this;
690  return tmp;
691  }
692 
693  __HOST_DEVICE__
694  HIP_vector_type& operator--() noexcept
695  {
696  return *this -= HIP_vector_type{1};
697  }
698  __HOST_DEVICE__
699  HIP_vector_type operator--(int) noexcept
700  {
701  auto tmp(*this);
702  --*this;
703  return tmp;
704  }
705 
706  __HOST_DEVICE__
707  HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
708  {
709  data += x.data;
710  return *this;
711  }
712  template<
713  typename U,
714  typename std::enable_if<
715  std::is_convertible<U, T>{}>::type* = nullptr>
716  __HOST_DEVICE__
717  HIP_vector_type& operator+=(U x) noexcept
718  {
719  return *this += HIP_vector_type{x};
720  }
721 
722  __HOST_DEVICE__
723  HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
724  {
725  data -= x.data;
726  return *this;
727  }
728  template<
729  typename U,
730  typename std::enable_if<
731  std::is_convertible<U, T>{}>::type* = nullptr>
732  __HOST_DEVICE__
733  HIP_vector_type& operator-=(U x) noexcept
734  {
735  return *this -= HIP_vector_type{x};
736  }
737 
738  __HOST_DEVICE__
739  HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
740  {
741  data *= x.data;
742  return *this;
743  }
744  template<
745  typename U,
746  typename std::enable_if<
747  std::is_convertible<U, T>{}>::type* = nullptr>
748  __HOST_DEVICE__
749  HIP_vector_type& operator*=(U x) noexcept
750  {
751  return *this *= HIP_vector_type{x};
752  }
753 
754  __HOST_DEVICE__
755  HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
756  {
757  data /= x.data;
758  return *this;
759  }
760  template<
761  typename U,
762  typename std::enable_if<
763  std::is_convertible<U, T>{}>::type* = nullptr>
764  __HOST_DEVICE__
765  HIP_vector_type& operator/=(U x) noexcept
766  {
767  return *this /= HIP_vector_type{x};
768  }
769 
770  template<
771  typename U = T,
772  typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
773  __HOST_DEVICE__
774  HIP_vector_type operator-() const noexcept
775  {
776  auto tmp(*this);
777  tmp.data = -tmp.data;
778  return tmp;
779  }
780 
781  template<
782  typename U = T,
783  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
784  __HOST_DEVICE__
785  HIP_vector_type operator~() const noexcept
786  {
787  HIP_vector_type r{*this};
788  r.data = ~r.data;
789  return r;
790  }
791 
792  template<
793  typename U = T,
794  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
795  __HOST_DEVICE__
796  HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
797  {
798  data %= x.data;
799  return *this;
800  }
801 
802  template<
803  typename U = T,
804  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
805  __HOST_DEVICE__
806  HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
807  {
808  data ^= x.data;
809  return *this;
810  }
811 
812  template<
813  typename U = T,
814  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
815  __HOST_DEVICE__
816  HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
817  {
818  data |= x.data;
819  return *this;
820  }
821 
822  template<
823  typename U = T,
824  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
825  __HOST_DEVICE__
826  HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
827  {
828  data &= x.data;
829  return *this;
830  }
831 
832  template<
833  typename U = T,
834  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
835  __HOST_DEVICE__
836  HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
837  {
838  data >>= x.data;
839  return *this;
840  }
841 
842  template<
843  typename U = T,
844  typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
845  __HOST_DEVICE__
846  HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
847  {
848  data <<= x.data;
849  return *this;
850  }
851  };
852 
853  template<typename T, unsigned int n>
854  __HOST_DEVICE__
855  inline
856  constexpr
857  HIP_vector_type<T, n> operator+(
858  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
859  {
860  return HIP_vector_type<T, n>{x} += y;
861  }
862  template<typename T, unsigned int n, typename U>
863  __HOST_DEVICE__
864  inline
865  constexpr
866  HIP_vector_type<T, n> operator+(
867  const HIP_vector_type<T, n>& x, U y) noexcept
868  {
869  return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
870  }
871  template<typename T, unsigned int n, typename U>
872  __HOST_DEVICE__
873  inline
874  constexpr
875  HIP_vector_type<T, n> operator+(
876  U x, const HIP_vector_type<T, n>& y) noexcept
877  {
878  return HIP_vector_type<T, n>{x} += y;
879  }
880 
881  template<typename T, unsigned int n>
882  __HOST_DEVICE__
883  inline
884  constexpr
885  HIP_vector_type<T, n> operator-(
886  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
887  {
888  return HIP_vector_type<T, n>{x} -= y;
889  }
890  template<typename T, unsigned int n, typename U>
891  __HOST_DEVICE__
892  inline
893  constexpr
894  HIP_vector_type<T, n> operator-(
895  const HIP_vector_type<T, n>& x, U y) noexcept
896  {
897  return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
898  }
899  template<typename T, unsigned int n, typename U>
900  __HOST_DEVICE__
901  inline
902  constexpr
903  HIP_vector_type<T, n> operator-(
904  U x, const HIP_vector_type<T, n>& y) noexcept
905  {
906  return HIP_vector_type<T, n>{x} -= y;
907  }
908 
909  template<typename T, unsigned int n>
910  __HOST_DEVICE__
911  inline
912  constexpr
913  HIP_vector_type<T, n> operator*(
914  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
915  {
916  return HIP_vector_type<T, n>{x} *= y;
917  }
918  template<typename T, unsigned int n, typename U>
919  __HOST_DEVICE__
920  inline
921  constexpr
922  HIP_vector_type<T, n> operator*(
923  const HIP_vector_type<T, n>& x, U y) noexcept
924  {
925  return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
926  }
927  template<typename T, unsigned int n, typename U>
928  __HOST_DEVICE__
929  inline
930  constexpr
931  HIP_vector_type<T, n> operator*(
932  U x, const HIP_vector_type<T, n>& y) noexcept
933  {
934  return HIP_vector_type<T, n>{x} *= y;
935  }
936 
937  template<typename T, unsigned int n>
938  __HOST_DEVICE__
939  inline
940  constexpr
941  HIP_vector_type<T, n> operator/(
942  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
943  {
944  return HIP_vector_type<T, n>{x} /= y;
945  }
946  template<typename T, unsigned int n, typename U>
947  __HOST_DEVICE__
948  inline
949  constexpr
950  HIP_vector_type<T, n> operator/(
951  const HIP_vector_type<T, n>& x, U y) noexcept
952  {
953  return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
954  }
955  template<typename T, unsigned int n, typename U>
956  __HOST_DEVICE__
957  inline
958  constexpr
959  HIP_vector_type<T, n> operator/(
960  U x, const HIP_vector_type<T, n>& y) noexcept
961  {
962  return HIP_vector_type<T, n>{x} /= y;
963  }
964 
965  template<typename V>
966  __HOST_DEVICE__
967  inline
968  constexpr
969  bool _hip_any_zero(const V& x, int n) noexcept
970  {
971  return
972  (n == -1) ? true : ((x[n] == 0) ? false : _hip_any_zero(x, n - 1));
973  }
974 
975  template<typename T, unsigned int n>
976  __HOST_DEVICE__
977  inline
978  constexpr
979  bool operator==(
980  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
981  {
982  return _hip_any_zero(x.data == y.data, n - 1);
983  }
984  template<typename T, unsigned int n, typename U>
985  __HOST_DEVICE__
986  inline
987  constexpr
988  bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
989  {
990  return x == HIP_vector_type<T, n>{y};
991  }
992  template<typename T, unsigned int n, typename U>
993  __HOST_DEVICE__
994  inline
995  constexpr
996  bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
997  {
998  return HIP_vector_type<T, n>{x} == y;
999  }
1000 
1001  template<typename T, unsigned int n>
1002  __HOST_DEVICE__
1003  inline
1004  constexpr
1005  bool operator!=(
1006  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1007  {
1008  return !(x == y);
1009  }
1010  template<typename T, unsigned int n, typename U>
1011  __HOST_DEVICE__
1012  inline
1013  constexpr
1014  bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
1015  {
1016  return !(x == y);
1017  }
1018  template<typename T, unsigned int n, typename U>
1019  __HOST_DEVICE__
1020  inline
1021  constexpr
1022  bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
1023  {
1024  return !(x == y);
1025  }
1026 
1027  template<
1028  typename T,
1029  unsigned int n,
1030  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1031  __HOST_DEVICE__
1032  inline
1033  constexpr
1034  HIP_vector_type<T, n> operator%(
1035  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1036  {
1037  return HIP_vector_type<T, n>{x} %= y;
1038  }
1039  template<
1040  typename T,
1041  unsigned int n,
1042  typename U,
1043  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1044  __HOST_DEVICE__
1045  inline
1046  constexpr
1047  HIP_vector_type<T, n> operator%(
1048  const HIP_vector_type<T, n>& x, U y) noexcept
1049  {
1050  return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
1051  }
1052  template<
1053  typename T,
1054  unsigned int n,
1055  typename U,
1056  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1057  __HOST_DEVICE__
1058  inline
1059  constexpr
1060  HIP_vector_type<T, n> operator%(
1061  U x, const HIP_vector_type<T, n>& y) noexcept
1062  {
1063  return HIP_vector_type<T, n>{x} %= y;
1064  }
1065 
1066  template<
1067  typename T,
1068  unsigned int n,
1069  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1070  __HOST_DEVICE__
1071  inline
1072  constexpr
1073  HIP_vector_type<T, n> operator^(
1074  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1075  {
1076  return HIP_vector_type<T, n>{x} ^= y;
1077  }
1078  template<
1079  typename T,
1080  unsigned int n,
1081  typename U,
1082  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1083  __HOST_DEVICE__
1084  inline
1085  constexpr
1086  HIP_vector_type<T, n> operator^(
1087  const HIP_vector_type<T, n>& x, U y) noexcept
1088  {
1089  return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
1090  }
1091  template<
1092  typename T,
1093  unsigned int n,
1094  typename U,
1095  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1096  __HOST_DEVICE__
1097  inline
1098  constexpr
1099  HIP_vector_type<T, n> operator^(
1100  U x, const HIP_vector_type<T, n>& y) noexcept
1101  {
1102  return HIP_vector_type<T, n>{x} ^= y;
1103  }
1104 
1105  template<
1106  typename T,
1107  unsigned int n,
1108  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1109  __HOST_DEVICE__
1110  inline
1111  constexpr
1112  HIP_vector_type<T, n> operator|(
1113  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1114  {
1115  return HIP_vector_type<T, n>{x} |= y;
1116  }
1117  template<
1118  typename T,
1119  unsigned int n,
1120  typename U,
1121  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1122  __HOST_DEVICE__
1123  inline
1124  constexpr
1125  HIP_vector_type<T, n> operator|(
1126  const HIP_vector_type<T, n>& x, U y) noexcept
1127  {
1128  return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
1129  }
1130  template<
1131  typename T,
1132  unsigned int n,
1133  typename U,
1134  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1135  __HOST_DEVICE__
1136  inline
1137  constexpr
1138  HIP_vector_type<T, n> operator|(
1139  U x, const HIP_vector_type<T, n>& y) noexcept
1140  {
1141  return HIP_vector_type<T, n>{x} |= y;
1142  }
1143 
1144  template<
1145  typename T,
1146  unsigned int n,
1147  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1148  __HOST_DEVICE__
1149  inline
1150  constexpr
1151  HIP_vector_type<T, n> operator&(
1152  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1153  {
1154  return HIP_vector_type<T, n>{x} &= y;
1155  }
1156  template<
1157  typename T,
1158  unsigned int n,
1159  typename U,
1160  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1161  __HOST_DEVICE__
1162  inline
1163  constexpr
1164  HIP_vector_type<T, n> operator&(
1165  const HIP_vector_type<T, n>& x, U y) noexcept
1166  {
1167  return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
1168  }
1169  template<
1170  typename T,
1171  unsigned int n,
1172  typename U,
1173  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1174  __HOST_DEVICE__
1175  inline
1176  constexpr
1177  HIP_vector_type<T, n> operator&(
1178  U x, const HIP_vector_type<T, n>& y) noexcept
1179  {
1180  return HIP_vector_type<T, n>{x} &= y;
1181  }
1182 
1183  template<
1184  typename T,
1185  unsigned int n,
1186  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1187  __HOST_DEVICE__
1188  inline
1189  constexpr
1190  HIP_vector_type<T, n> operator>>(
1191  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1192  {
1193  return HIP_vector_type<T, n>{x} >>= y;
1194  }
1195  template<
1196  typename T,
1197  unsigned int n,
1198  typename U,
1199  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1200  __HOST_DEVICE__
1201  inline
1202  constexpr
1203  HIP_vector_type<T, n> operator>>(
1204  const HIP_vector_type<T, n>& x, U y) noexcept
1205  {
1206  return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
1207  }
1208  template<
1209  typename T,
1210  unsigned int n,
1211  typename U,
1212  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1213  __HOST_DEVICE__
1214  inline
1215  constexpr
1216  HIP_vector_type<T, n> operator>>(
1217  U x, const HIP_vector_type<T, n>& y) noexcept
1218  {
1219  return HIP_vector_type<T, n>{x} >>= y;
1220  }
1221 
1222  template<
1223  typename T,
1224  unsigned int n,
1225  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1226  __HOST_DEVICE__
1227  inline
1228  constexpr
1229  HIP_vector_type<T, n> operator<<(
1230  const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
1231  {
1232  return HIP_vector_type<T, n>{x} <<= y;
1233  }
1234  template<
1235  typename T,
1236  unsigned int n,
1237  typename U,
1238  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1239  __HOST_DEVICE__
1240  inline
1241  constexpr
1242  HIP_vector_type<T, n> operator<<(
1243  const HIP_vector_type<T, n>& x, U y) noexcept
1244  {
1245  return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
1246  }
1247  template<
1248  typename T,
1249  unsigned int n,
1250  typename U,
1251  typename std::enable_if<std::is_arithmetic<U>::value>::type,
1252  typename std::enable_if<std::is_integral<T>{}>* = nullptr>
1253  __HOST_DEVICE__
1254  inline
1255  constexpr
1256  HIP_vector_type<T, n> operator<<(
1257  U x, const HIP_vector_type<T, n>& y) noexcept
1258  {
1259  return HIP_vector_type<T, n>{x} <<= y;
1260  }
1261 
1262  #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
1263  using CUDA_name##1 = HIP_vector_type<T, 1>;\
1264  using CUDA_name##2 = HIP_vector_type<T, 2>;\
1265  using CUDA_name##3 = HIP_vector_type<T, 3>;\
1266  using CUDA_name##4 = HIP_vector_type<T, 4>;
1267 #else
1268  #define __MAKE_VECTOR_TYPE__(CUDA_name, T) \
1269  typedef struct {\
1270  T x;\
1271  } CUDA_name##1;\
1272  typedef struct {\
1273  T x;\
1274  T y;\
1275  } CUDA_name##2;\
1276  typedef struct {\
1277  T x;\
1278  T y;\
1279  T z;\
1280  } CUDA_name##3;\
1281  typedef struct {\
1282  T x;\
1283  T y;\
1284  T z;\
1285  T w;\
1286  } CUDA_name##4;
1287 #endif
1288 
1289 __MAKE_VECTOR_TYPE__(uchar, unsigned char);
1290 __MAKE_VECTOR_TYPE__(char, char);
1291 __MAKE_VECTOR_TYPE__(ushort, unsigned short);
1292 __MAKE_VECTOR_TYPE__(short, short);
1293 __MAKE_VECTOR_TYPE__(uint, unsigned int);
1294 __MAKE_VECTOR_TYPE__(int, int);
1295 __MAKE_VECTOR_TYPE__(ulong, unsigned long);
1296 __MAKE_VECTOR_TYPE__(long, long);
1297 __MAKE_VECTOR_TYPE__(ulonglong, unsigned long long);
1298 __MAKE_VECTOR_TYPE__(longlong, long long);
1299 __MAKE_VECTOR_TYPE__(float, float);
1300 __MAKE_VECTOR_TYPE__(double, double);
1301 
1302 #ifdef __cplusplus
1303 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
1304  static inline __HOST_DEVICE__ \
1305  type make_##type(comp x) { type r{x}; return r; }
1306 
1307 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
1308  static inline __HOST_DEVICE__ \
1309  type make_##type(comp x, comp y) { type r{x, y}; return r; }
1310 
1311 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
1312  static inline __HOST_DEVICE__ \
1313  type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; }
1314 
1315 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
1316  static inline __HOST_DEVICE__ \
1317  type make_##type(comp x, comp y, comp z, comp w) { \
1318  type r{x, y, z, w}; \
1319  return r; \
1320  }
1321 #else
1322  #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
1323  static inline __HOST_DEVICE__ \
1324  type make_##type(comp x) { type r; r.x =x; return r; }
1325 
1326  #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
1327  static inline __HOST_DEVICE__ \
1328  type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; }
1329 
1330  #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
1331  static inline __HOST_DEVICE__ \
1332  type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; }
1333 
1334  #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
1335  static inline __HOST_DEVICE__ \
1336  type make_##type(comp x, comp y, comp z, comp w) { \
1337  type r; r.x=x; r.y=y; r.z=z; r.w=w; \
1338  return r; \
1339  }
1340 #endif
1341 
1342 DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1);
1343 DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2);
1344 DECLOP_MAKE_THREE_COMPONENT(unsigned char, uchar3);
1345 DECLOP_MAKE_FOUR_COMPONENT(unsigned char, uchar4);
1346 
1347 DECLOP_MAKE_ONE_COMPONENT(signed char, char1);
1348 DECLOP_MAKE_TWO_COMPONENT(signed char, char2);
1349 DECLOP_MAKE_THREE_COMPONENT(signed char, char3);
1350 DECLOP_MAKE_FOUR_COMPONENT(signed char, char4);
1351 
1352 DECLOP_MAKE_ONE_COMPONENT(unsigned short, ushort1);
1353 DECLOP_MAKE_TWO_COMPONENT(unsigned short, ushort2);
1354 DECLOP_MAKE_THREE_COMPONENT(unsigned short, ushort3);
1355 DECLOP_MAKE_FOUR_COMPONENT(unsigned short, ushort4);
1356 
1357 DECLOP_MAKE_ONE_COMPONENT(signed short, short1);
1358 DECLOP_MAKE_TWO_COMPONENT(signed short, short2);
1359 DECLOP_MAKE_THREE_COMPONENT(signed short, short3);
1360 DECLOP_MAKE_FOUR_COMPONENT(signed short, short4);
1361 
1362 DECLOP_MAKE_ONE_COMPONENT(unsigned int, uint1);
1363 DECLOP_MAKE_TWO_COMPONENT(unsigned int, uint2);
1364 DECLOP_MAKE_THREE_COMPONENT(unsigned int, uint3);
1365 DECLOP_MAKE_FOUR_COMPONENT(unsigned int, uint4);
1366 
1367 DECLOP_MAKE_ONE_COMPONENT(signed int, int1);
1368 DECLOP_MAKE_TWO_COMPONENT(signed int, int2);
1369 DECLOP_MAKE_THREE_COMPONENT(signed int, int3);
1370 DECLOP_MAKE_FOUR_COMPONENT(signed int, int4);
1371 
1372 DECLOP_MAKE_ONE_COMPONENT(float, float1);
1373 DECLOP_MAKE_TWO_COMPONENT(float, float2);
1374 DECLOP_MAKE_THREE_COMPONENT(float, float3);
1375 DECLOP_MAKE_FOUR_COMPONENT(float, float4);
1376 
1377 DECLOP_MAKE_ONE_COMPONENT(double, double1);
1378 DECLOP_MAKE_TWO_COMPONENT(double, double2);
1379 DECLOP_MAKE_THREE_COMPONENT(double, double3);
1380 DECLOP_MAKE_FOUR_COMPONENT(double, double4);
1381 
1382 DECLOP_MAKE_ONE_COMPONENT(unsigned long, ulong1);
1383 DECLOP_MAKE_TWO_COMPONENT(unsigned long, ulong2);
1384 DECLOP_MAKE_THREE_COMPONENT(unsigned long, ulong3);
1385 DECLOP_MAKE_FOUR_COMPONENT(unsigned long, ulong4);
1386 
1387 DECLOP_MAKE_ONE_COMPONENT(signed long, long1);
1388 DECLOP_MAKE_TWO_COMPONENT(signed long, long2);
1389 DECLOP_MAKE_THREE_COMPONENT(signed long, long3);
1390 DECLOP_MAKE_FOUR_COMPONENT(signed long, long4);
1391 
1392 DECLOP_MAKE_ONE_COMPONENT(unsigned long long, ulonglong1);
1393 DECLOP_MAKE_TWO_COMPONENT(unsigned long long, ulonglong2);
1394 DECLOP_MAKE_THREE_COMPONENT(unsigned long long, ulonglong3);
1395 DECLOP_MAKE_FOUR_COMPONENT(unsigned long long, ulonglong4);
1396 
1397 DECLOP_MAKE_ONE_COMPONENT(signed long long, longlong1);
1398 DECLOP_MAKE_TWO_COMPONENT(signed long long, longlong2);
1399 DECLOP_MAKE_THREE_COMPONENT(signed long long, longlong3);
1400 DECLOP_MAKE_FOUR_COMPONENT(signed long long, longlong4);
1401 #else // !defined(__has_attribute)
1402 
1403 #if defined(_MSC_VER)
1404 #include <mmintrin.h>
1405 #include <xmmintrin.h>
1406 #include <emmintrin.h>
1407 #include <immintrin.h>
1408 
1409 typedef union { char data; } char1;
1410 typedef union { char data[2]; } char2;
1411 typedef union { char data[4]; } char4;
1412 typedef union { char4 data; } char3;
1413 typedef union { __m64 data; } char8;
1414 typedef union { __m128i data; } char16;
1415 
1416 typedef union { unsigned char data; } uchar1;
1417 typedef union { unsigned char data[2]; } uchar2;
1418 typedef union { unsigned char data[4]; } uchar4;
1419 typedef union { uchar4 data; } uchar3;
1420 typedef union { __m64 data; } uchar8;
1421 typedef union { __m128i data; } uchar16;
1422 
1423 typedef union { short data; } short1;
1424 typedef union { short data[2]; } short2;
1425 typedef union { __m64 data; } short4;
1426 typedef union { short4 data; } short3;
1427 typedef union { __m128i data; } short8;
1428 typedef union { __m128i data[2]; } short16;
1429 
1430 typedef union { unsigned short data; } ushort1;
1431 typedef union { unsigned short data[2]; } ushort2;
1432 typedef union { __m64 data; } ushort4;
1433 typedef union { ushort4 data; } ushort3;
1434 typedef union { __m128i data; } ushort8;
1435 typedef union { __m128i data[2]; } ushort16;
1436 
1437 typedef union { int data; } int1;
1438 typedef union { __m64 data; } int2;
1439 typedef union { __m128i data; } int4;
1440 typedef union { int4 data; } int3;
1441 typedef union { __m128i data[2]; } int8;
1442 typedef union { __m128i data[4];} int16;
1443 
1444 typedef union { unsigned int data; } uint1;
1445 typedef union { __m64 data; } uint2;
1446 typedef union { __m128i data; } uint4;
1447 typedef union { uint4 data; } uint3;
1448 typedef union { __m128i data[2]; } uint8;
1449 typedef union { __m128i data[4]; } uint16;
1450 
1451 #if !defined(_WIN64)
1452 typedef union { int data; } long1;
1453 typedef union { __m64 data; } long2;
1454 typedef union { __m128i data; } long4;
1455 typedef union { long4 data; } long3;
1456 typedef union { __m128i data[2]; } long8;
1457 typedef union { __m128i data[4]; } long16;
1458 
1459 typedef union { unsigned int data; } ulong1;
1460 typedef union { __m64 data; } ulong2;
1461 typedef union { __m128i data; } ulong4;
1462 typedef union { ulong4 data; } ulong3;
1463 typedef union { __m128i data[2]; } ulong8;
1464 typedef union { __m128i data[4]; } ulong16;
1465 #else // defined(_WIN64)
1466 typedef union { __m64 data; } long1;
1467 typedef union { __m128i data; } long2;
1468 typedef union { __m128i data[2]; } long4;
1469 typedef union { long4 data; } long3;
1470 typedef union { __m128i data[4]; } long8;
1471 typedef union { __m128i data[8]; } long16;
1472 
1473 typedef union { __m64 data; } ulong1;
1474 typedef union { __m128i data; } ulong2;
1475 typedef union { __m128i data[2]; } ulong4;
1476 typedef union { ulong4 data; } ulong3;
1477 typedef union { __m128i data[4]; } ulong8;
1478 typedef union { __m128i data[8]; } ulong16;
1479 #endif // defined(_WIN64)
1480 
1481 typedef union { __m64 data; } longlong1;
1482 typedef union { __m128i data; } longlong2;
1483 typedef union { __m128i data[2]; } longlong4;
1484 typedef union { longlong4 data; } longlong3;
1485 typedef union { __m128i data[4]; } longlong8;
1486 typedef union { __m128i data[8]; } longlong16;
1487 
1488 typedef union { __m64 data; } ulonglong1;
1489 typedef union { __m128i data; } ulonglong2;
1490 typedef union { __m128i data[2]; } ulonglong4;
1491 typedef union { ulonglong4 data; } ulonglong3;
1492 typedef union { __m128i data[4]; } ulonglong8;
1493 typedef union { __m128i data[8]; } ulonglong16;
1494 
1495 typedef union { float data; } float1;
1496 typedef union { __m64 data; } float2;
1497 typedef union { __m128 data; } float4;
1498 typedef union { float4 data; } float3;
1499 typedef union { __m256 data; } float8;
1500 typedef union { __m256 data[2]; } float16;
1501 
1502 typedef union { double data; } double1;
1503 typedef union { __m128d data; } double2;
1504 typedef union { __m256d data; } double4;
1505 typedef union { double4 data; } double3;
1506 typedef union { __m256d data[2]; } double8;
1507 typedef union { __m256d data[4]; } double16;
1508 
1509 #else // !defined(_MSC_VER)
1510 
1511 typedef union { char data; } char1;
1512 typedef union { char data[2]; } char2;
1513 typedef union { char data[4]; } char4;
1514 typedef union { char data[8]; } char8;
1515 typedef union { char data[16]; } char16;
1516 typedef union { char4 data; } char3;
1517 
1518 typedef union { unsigned char data; } uchar1;
1519 typedef union { unsigned char data[2]; } uchar2;
1520 typedef union { unsigned char data[4]; } uchar4;
1521 typedef union { unsigned char data[8]; } uchar8;
1522 typedef union { unsigned char data[16]; } uchar16;
1523 typedef union { uchar4 data; } uchar3;
1524 
1525 typedef union { short data; } short1;
1526 typedef union { short data[2]; } short2;
1527 typedef union { short data[4]; } short4;
1528 typedef union { short data[8]; } short8;
1529 typedef union { short data[16]; } short16;
1530 typedef union { short4 data; } short3;
1531 
1532 typedef union { unsigned short data; } ushort1;
1533 typedef union { unsigned short data[2]; } ushort2;
1534 typedef union { unsigned short data[4]; } ushort4;
1535 typedef union { unsigned short data[8]; } ushort8;
1536 typedef union { unsigned short data[16]; } ushort16;
1537 typedef union { ushort4 data; } ushort3;
1538 
1539 typedef union { int data; } int1;
1540 typedef union { int data[2]; } int2;
1541 typedef union { int data[4]; } int4;
1542 typedef union { int data[8]; } int8;
1543 typedef union { int data[16]; } int16;
1544 typedef union { int4 data; } int3;
1545 
1546 typedef union { unsigned int data; } uint1;
1547 typedef union { unsigned int data[2]; } uint2;
1548 typedef union { unsigned int data[4]; } uint4;
1549 typedef union { unsigned int data[8]; } uint8;
1550 typedef union { unsigned int data[16]; } uint16;
1551 typedef union { uint4 data; } uint3;
1552 
1553 typedef union { long data; } long1;
1554 typedef union { long data[2]; } long2;
1555 typedef union { long data[4]; } long4;
1556 typedef union { long data[8]; } long8;
1557 typedef union { long data[16]; } long16;
1558 typedef union { long4 data; } long3;
1559 
1560 typedef union { unsigned long data; } ulong1;
1561 typedef union { unsigned long data[2]; } ulong2;
1562 typedef union { unsigned long data[4]; } ulong4;
1563 typedef union { unsigned long data[8]; } ulong8;
1564 typedef union { unsigned long data[16]; } ulong16;
1565 typedef union { ulong4 data; } ulong3;
1566 
1567 typedef union { long long data; } longlong1;
1568 typedef union { long long data[2]; } longlong2;
1569 typedef union { long long data[4]; } longlong4;
1570 typedef union { long long data[8]; } longlong8;
1571 typedef union { long long data[16]; } longlong16;
1572 typedef union { longlong4 data; } longlong3;
1573 
1574 typedef union { unsigned long long data; } ulonglong1;
1575 typedef union { unsigned long long data[2]; } ulonglong2;
1576 typedef union { unsigned long long data[4]; } ulonglong4;
1577 typedef union { unsigned long long data[8]; } ulonglong8;
1578 typedef union { unsigned long long data[16]; } ulonglong16;
1579 typedef union { ulonglong4 data; } ulonglong3;
1580 
1581 typedef union { float data; } float1;
1582 typedef union { float data[2]; } float2;
1583 typedef union { float data[4]; } float4;
1584 typedef union { float data[8]; } float8;
1585 typedef union { float data[16]; } float16;
1586 typedef union { float4 data; } float3;
1587 
1588 typedef union { double data; } double1;
1589 typedef union { double data[2]; } double2;
1590 typedef union { double data[4]; } double4;
1591 typedef union { double data[8]; } double8;
1592 typedef union { double data[16]; } double16;
1593 typedef union { double4 data; } double3;
1594 
1595 #endif // defined(_MSC_VER)
1596 #endif // defined(__has_attribute)
1597 #endif
uint2
Definition: hip_vector_types.h:1547
longlong2
Definition: hip_vector_types.h:1568
int4
Definition: hip_vector_types.h:1541
double8
Definition: hip_vector_types.h:1591
long2
Definition: hip_vector_types.h:1554
ulong16
Definition: hip_vector_types.h:1564
uchar2
Definition: hip_vector_types.h:1519
ulong4
Definition: hip_vector_types.h:1562
char4
Definition: hip_vector_types.h:1513
ushort1
Definition: hip_vector_types.h:1532
char2
Definition: hip_vector_types.h:1512
int8
Definition: hip_vector_types.h:1542
ulong1
Definition: hip_vector_types.h:1560
double2
Definition: hip_vector_types.h:1589
uint3
Definition: hip_vector_types.h:1551
long3
Definition: hip_vector_types.h:1558
ulong8
Definition: hip_vector_types.h:1563
long16
Definition: hip_vector_types.h:1557
uint1
Definition: hip_vector_types.h:1546
int3
Definition: hip_vector_types.h:1544
long1
Definition: hip_vector_types.h:1553
uint16
Definition: hip_vector_types.h:1550
float4
Definition: hip_vector_types.h:1583
double16
Definition: hip_vector_types.h:1592
char3
Definition: hip_vector_types.h:1516
short3
Definition: hip_vector_types.h:1530
char1
Definition: hip_vector_types.h:1511
longlong3
Definition: hip_vector_types.h:1572
ulong2
Definition: hip_vector_types.h:1561
ulonglong8
Definition: hip_vector_types.h:1577
float8
Definition: hip_vector_types.h:1584
int2
Definition: hip_vector_types.h:1540
host_defines.h
TODO-doc.
float3
Definition: hip_vector_types.h:1586
ulonglong16
Definition: hip_vector_types.h:1578
uchar3
Definition: hip_vector_types.h:1523
ulonglong4
Definition: hip_vector_types.h:1576
float2
Definition: hip_vector_types.h:1582
ushort16
Definition: hip_vector_types.h:1536
short4
Definition: hip_vector_types.h:1527
longlong4
Definition: hip_vector_types.h:1569
uchar16
Definition: hip_vector_types.h:1522
ushort8
Definition: hip_vector_types.h:1535
short1
Definition: hip_vector_types.h:1525
int1
Definition: hip_vector_types.h:1539
double3
Definition: hip_vector_types.h:1593
char16
Definition: hip_vector_types.h:1515
ulonglong2
Definition: hip_vector_types.h:1575
short8
Definition: hip_vector_types.h:1528
longlong8
Definition: hip_vector_types.h:1570
uchar4
Definition: hip_vector_types.h:1520
ulonglong3
Definition: hip_vector_types.h:1579
ushort4
Definition: hip_vector_types.h:1534
float16
Definition: hip_vector_types.h:1585
float1
Definition: hip_vector_types.h:1581
short16
Definition: hip_vector_types.h:1529
longlong1
Definition: hip_vector_types.h:1567
uchar1
Definition: hip_vector_types.h:1518
uint8
Definition: hip_vector_types.h:1549
short2
Definition: hip_vector_types.h:1526
long8
Definition: hip_vector_types.h:1556
ulong3
Definition: hip_vector_types.h:1565
uchar8
Definition: hip_vector_types.h:1521
double4
Definition: hip_vector_types.h:1590
longlong16
Definition: hip_vector_types.h:1571
ulonglong1
Definition: hip_vector_types.h:1574
ushort2
Definition: hip_vector_types.h:1533
double1
Definition: hip_vector_types.h:1588
ushort3
Definition: hip_vector_types.h:1537
char8
Definition: hip_vector_types.h:1514
uint4
Definition: hip_vector_types.h:1548
int16
Definition: hip_vector_types.h:1543
long4
Definition: hip_vector_types.h:1555