HIP: Heterogenous-computing Interface for Portability
hip_fp16.h
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 
23 #pragma once
25 #include <assert.h>
26 #if defined(__cplusplus)
27  #include <algorithm>
28  #include <type_traits>
29  #include <utility>
30 #endif
31 
32 #if defined(__clang__) && (__clang_major__ > 5)
33  typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2)));
34 
35  struct __half_raw {
36  union {
37  static_assert(sizeof(_Float16) == sizeof(unsigned short), "");
38 
39  _Float16 data;
40  unsigned short x;
41  };
42  };
43 
44  struct __half2_raw {
45  union {
46  static_assert(sizeof(_Float16_2) == sizeof(unsigned short[2]), "");
47 
48  _Float16_2 data;
49  struct {
50  unsigned short x;
51  unsigned short y;
52  };
53  };
54  };
55 
56  #if defined(__cplusplus)
57  #include "hip_fp16_math_fwd.h"
58  #include "hip_vector_types.h"
59  #include "host_defines.h"
60 
61  namespace std
62  {
63  template<> struct is_floating_point<_Float16> : std::true_type {};
64  }
65 
66  template<bool cond, typename T = void>
67  using Enable_if_t = typename std::enable_if<cond, T>::type;
68 
69  // BEGIN STRUCT __HALF
70  struct __half {
71  protected:
72  union {
73  static_assert(sizeof(_Float16) == sizeof(unsigned short), "");
74 
75  _Float16 data;
76  unsigned short __x;
77  };
78  public:
79  // CREATORS
80  __host__ __device__
81  __half() = default;
82  __host__ __device__
83  __half(const __half_raw& x) : data{x.data} {}
84  #if !defined(__HIP_NO_HALF_CONVERSIONS__)
85  __host__ __device__
86  __half(decltype(data) x) : data{x} {}
87  template<
88  typename T,
89  Enable_if_t<std::is_floating_point<T>{}>* = nullptr>
90  __host__ __device__
91  __half(T x) : data{static_cast<_Float16>(x)} {}
92  #endif
93  __host__ __device__
94  __half(const __half&) = default;
95  __host__ __device__
96  __half(__half&&) = default;
97  __host__ __device__
98  ~__half() = default;
99 
100  // CREATORS - DEVICE ONLY
101  #if !defined(__HIP_NO_HALF_CONVERSIONS__)
102  template<
103  typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
104  __device__
105  __half(T x) : data{static_cast<_Float16>(x)} {}
106  #endif
107 
108  // MANIPULATORS
109  __host__ __device__
110  __half& operator=(const __half&) = default;
111  __host__ __device__
112  __half& operator=(__half&&) = default;
113  __host__ __device__
114  __half& operator=(const __half_raw& x)
115  {
116  data = x.data;
117  return *this;
118  }
119  __host__ __device__
120  volatile __half& operator=(const __half_raw& x) volatile
121  {
122  data = x.data;
123  return *this;
124  }
125  volatile __half& operator=(const volatile __half_raw& x) volatile
126  {
127  data = x.data;
128  return *this;
129  }
130  __half& operator=(__half_raw&& x)
131  {
132  data = x.data;
133  return *this;
134  }
135  volatile __half& operator=(__half_raw&& x) volatile
136  {
137  data = x.data;
138  return *this;
139  }
140  volatile __half& operator=(volatile __half_raw&& x) volatile
141  {
142  data = x.data;
143  return *this;
144  }
145  #if !defined(__HIP_NO_HALF_CONVERSIONS__)
146  template<
147  typename T,
148  Enable_if_t<std::is_floating_point<T>{}>* = nullptr>
149  __host__ __device__
150  __half& operator=(T x)
151  {
152  data = static_cast<_Float16>(x);
153  return *this;
154  }
155  #endif
156 
157  // MANIPULATORS - DEVICE ONLY
158  #if !defined(__HIP_NO_HALF_CONVERSIONS__)
159  template<
160  typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
161  __device__
162  __half& operator=(T x)
163  {
164  data = static_cast<_Float16>(x);
165  return *this;
166  }
167  #endif
168 
169  #if !defined(__HIP_NO_HALF_OPERATORS__)
170  __device__
171  __half& operator+=(const __half& x)
172  {
173  data += x.data;
174  return *this;
175  }
176  __device__
177  __half& operator-=(const __half& x)
178  {
179  data -= x.data;
180  return *this;
181  }
182  __device__
183  __half& operator*=(const __half& x)
184  {
185  data *= x.data;
186  return *this;
187  }
188  __device__
189  __half& operator/=(const __half& x)
190  {
191  data /= x.data;
192  return *this;
193  }
194  __device__
195  __half& operator++() { ++data; return *this; }
196  __device__
197  __half operator++(int)
198  {
199  __half tmp{*this};
200  ++*this;
201  return tmp;
202  }
203  __device__
204  __half& operator--() { --data; return *this; }
205  __device__
206  __half operator--(int)
207  {
208  __half tmp{*this};
209  --*this;
210  return tmp;
211  }
212  #endif
213 
214  // ACCESSORS
215  #if !defined(__HIP_NO_HALF_CONVERSIONS__)
216  template<
217  typename T,
218  Enable_if_t<
219  std::is_floating_point<T>{} &&
220  !std::is_same<T, double>{}>* = nullptr>
221  operator T() const { return data; }
222  #endif
223  __host__ __device__
224  operator __half_raw() const { return __half_raw{data}; }
225  __host__ __device__
226  operator volatile __half_raw() const volatile
227  {
228  return __half_raw{data};
229  }
230 
231  // ACCESSORS - DEVICE ONLY
232  #if !defined(__HIP_NO_HALF_CONVERSIONS__)
233  template<
234  typename T, Enable_if_t<std::is_integral<T>{}>* = nullptr>
235  __device__
236  operator T() const { return data; }
237  #endif
238 
239  #if !defined(__HIP_NO_HALF_OPERATORS__)
240  __device__
241  __half operator+() const { return *this; }
242  __device__
243  __half operator-() const
244  {
245  __half tmp{*this};
246  tmp.data = -tmp.data;
247  return tmp;
248  }
249  #endif
250 
251  // FRIENDS
252  #if !defined(__HIP_NO_HALF_OPERATORS__)
253  friend
254  inline
255  __device__
256  __half operator+(const __half& x, const __half& y)
257  {
258  return __half{x} += y;
259  }
260  friend
261  inline
262  __device__
263  __half operator-(const __half& x, const __half& y)
264  {
265  return __half{x} -= y;
266  }
267  friend
268  inline
269  __device__
270  __half operator*(const __half& x, const __half& y)
271  {
272  return __half{x} *= y;
273  }
274  friend
275  inline
276  __device__
277  __half operator/(const __half& x, const __half& y)
278  {
279  return __half{x} /= y;
280  }
281  friend
282  inline
283  __device__
284  bool operator==(const __half& x, const __half& y)
285  {
286  return x.data == y.data;
287  }
288  friend
289  inline
290  __device__
291  bool operator!=(const __half& x, const __half& y)
292  {
293  return !(x == y);
294  }
295  friend
296  inline
297  __device__
298  bool operator<(const __half& x, const __half& y)
299  {
300  return x.data < y.data;
301  }
302  friend
303  inline
304  __device__
305  bool operator>(const __half& x, const __half& y)
306  {
307  return y.data < x.data;
308  }
309  friend
310  inline
311  __device__
312  bool operator<=(const __half& x, const __half& y)
313  {
314  return !(y < x);
315  }
316  friend
317  inline
318  __device__
319  bool operator>=(const __half& x, const __half& y)
320  {
321  return !(x < y);
322  }
323  #endif // !defined(__HIP_NO_HALF_OPERATORS__)
324  };
325  // END STRUCT __HALF
326 
327  // BEGIN STRUCT __HALF2
328  struct __half2 {
329  protected:
330  union {
331  static_assert(
332  sizeof(_Float16_2) == sizeof(unsigned short[2]), "");
333 
334  _Float16_2 data;
335  struct {
336  unsigned short x;
337  unsigned short y;
338  };
339  };
340  public:
341  // CREATORS
342  __host__ __device__
343  __half2() = default;
344  __host__ __device__
345  __half2(const __half2_raw& x) : data{x.data} {}
346  __host__ __device__
347  __half2(decltype(data) x) : data{x} {}
348  __host__ __device__
349  __half2(const __half& x, const __half& y)
350  :
351  data{
352  static_cast<__half_raw>(x).data,
353  static_cast<__half_raw>(y).data}
354  {}
355  __host__ __device__
356  __half2(const __half2&) = default;
357  __host__ __device__
358  __half2(__half2&&) = default;
359  __host__ __device__
360  ~__half2() = default;
361 
362  // MANIPULATORS
363  __host__ __device__
364  __half2& operator=(const __half2&) = default;
365  __host__ __device__
366  __half2& operator=(__half2&&) = default;
367  __host__ __device__
368  __half2& operator=(const __half2_raw& x)
369  {
370  data = x.data;
371  return *this;
372  }
373 
374  // MANIPULATORS - DEVICE ONLY
375  #if !defined(__HIP_NO_HALF_OPERATORS__)
376  __device__
377  __half2& operator+=(const __half2& x)
378  {
379  data += x.data;
380  return *this;
381  }
382  __device__
383  __half2& operator-=(const __half2& x)
384  {
385  data -= x.data;
386  return *this;
387  }
388  __device__
389  __half2& operator*=(const __half2& x)
390  {
391  data *= x.data;
392  return *this;
393  }
394  __device__
395  __half2& operator/=(const __half2& x)
396  {
397  data /= x.data;
398  return *this;
399  }
400  __device__
401  __half2& operator++() { return *this += _Float16_2{1, 1}; }
402  __device__
403  __half2 operator++(int)
404  {
405  __half2 tmp{*this};
406  ++*this;
407  return tmp;
408  }
409  __device__
410  __half2& operator--() { return *this -= _Float16_2{1, 1}; }
411  __device__
412  __half2 operator--(int)
413  {
414  __half2 tmp{*this};
415  --*this;
416  return tmp;
417  }
418  #endif
419 
420  // ACCESSORS
421  __host__ __device__
422  operator decltype(data)() const { return data; }
423  __host__ __device__
424  operator __half2_raw() const { return __half2_raw{data}; }
425 
426  // ACCESSORS - DEVICE ONLY
427  #if !defined(__HIP_NO_HALF_OPERATORS__)
428  __device__
429  __half2 operator+() const { return *this; }
430  __device__
431  __half2 operator-() const
432  {
433  __half2 tmp{*this};
434  tmp.data = -tmp.data;
435  return tmp;
436  }
437  #endif
438 
439  // FRIENDS
440  #if !defined(__HIP_NO_HALF_OPERATORS__)
441  friend
442  inline
443  __device__
444  __half2 operator+(const __half2& x, const __half2& y)
445  {
446  return __half2{x} += y;
447  }
448  friend
449  inline
450  __device__
451  __half2 operator-(const __half2& x, const __half2& y)
452  {
453  return __half2{x} -= y;
454  }
455  friend
456  inline
457  __device__
458  __half2 operator*(const __half2& x, const __half2& y)
459  {
460  return __half2{x} *= y;
461  }
462  friend
463  inline
464  __device__
465  __half2 operator/(const __half2& x, const __half2& y)
466  {
467  return __half2{x} /= y;
468  }
469  friend
470  inline
471  __device__
472  bool operator==(const __half2& x, const __half2& y)
473  {
474  auto r = x.data == y.data;
475  return r.x != 0 && r.y != 0;
476  }
477  friend
478  inline
479  __device__
480  bool operator!=(const __half2& x, const __half2& y)
481  {
482  return !(x == y);
483  }
484  friend
485  inline
486  __device__
487  bool operator<(const __half2& x, const __half2& y)
488  {
489  auto r = x.data < y.data;
490  return r.x != 0 && r.y != 0;
491  }
492  friend
493  inline
494  __device__
495  bool operator>(const __half2& x, const __half2& y)
496  {
497  return y < x;
498  }
499  friend
500  inline
501  __device__
502  bool operator<=(const __half2& x, const __half2& y)
503  {
504  return !(y < x);
505  }
506  friend
507  inline
508  __device__
509  bool operator>=(const __half2& x, const __half2& y)
510  {
511  return !(x < y);
512  }
513  #endif // !defined(__HIP_NO_HALF_OPERATORS__)
514  };
515  // END STRUCT __HALF2
516 
517  namespace
518  {
519  inline
520  __host__ __device__
521  __half2 make_half2(__half x, __half y)
522  {
523  return __half2{x, y};
524  }
525 
526  inline
527  __device__
528  __half __low2half(__half2 x)
529  {
530  return __half{__half_raw{static_cast<__half2_raw>(x).data.x}};
531  }
532 
533  inline
534  __device__
535  __half __high2half(__half2 x)
536  {
537  return __half{__half_raw{static_cast<__half2_raw>(x).data.y}};
538  }
539 
540  inline
541  __device__
542  __half2 __half2half2(__half x)
543  {
544  return __half2{x, x};
545  }
546 
547  inline
548  __device__
549  __half2 __halves2half2(__half x, __half y)
550  {
551  return __half2{x, y};
552  }
553 
554  inline
555  __device__
556  __half2 __low2half2(__half2 x)
557  {
558  return __half2{
559  _Float16_2{
560  static_cast<__half2_raw>(x).data.x,
561  static_cast<__half2_raw>(x).data.x}};
562  }
563 
564  inline
565  __device__
566  __half2 __high2half2(__half2 x)
567  {
568  return __half2_raw{
569  _Float16_2{
570  static_cast<__half2_raw>(x).data.y,
571  static_cast<__half2_raw>(x).data.y}};
572  }
573 
574  inline
575  __device__
576  __half2 __lows2half2(__half2 x, __half2 y)
577  {
578  return __half2_raw{
579  _Float16_2{
580  static_cast<__half2_raw>(x).data.x,
581  static_cast<__half2_raw>(y).data.x}};
582  }
583 
584  inline
585  __device__
586  __half2 __highs2half2(__half2 x, __half2 y)
587  {
588  return __half2_raw{
589  _Float16_2{
590  static_cast<__half2_raw>(x).data.y,
591  static_cast<__half2_raw>(y).data.y}};
592  }
593 
594  inline
595  __device__
596  __half2 __lowhigh2highlow(__half2 x)
597  {
598  return __half2_raw{
599  _Float16_2{
600  static_cast<__half2_raw>(x).data.y,
601  static_cast<__half2_raw>(x).data.x}};
602  }
603 
604  // Bitcasts
605  inline
606  __device__
607  short __half_as_short(__half x)
608  {
609  return static_cast<__half_raw>(x).x;
610  }
611 
612  inline
613  __device__
614  unsigned short __half_as_ushort(__half x)
615  {
616  return static_cast<__half_raw>(x).x;
617  }
618 
619  inline
620  __device__
621  __half __short_as_half(short x)
622  {
623  __half_raw r; r.x = x;
624  return r;
625  }
626 
627  inline
628  __device__
629  __half __ushort_as_half(unsigned short x)
630  {
631  __half_raw r; r.x = x;
632  return r;
633  }
634 
635  // TODO: rounding behaviour is not correct.
636  // float -> half | half2
637  inline
638  __device__ __host__
639  __half __float2half(float x)
640  {
641  return __half_raw{static_cast<_Float16>(x)};
642  }
643  inline
644  __device__ __host__
645  __half __float2half_rn(float x)
646  {
647  return __half_raw{static_cast<_Float16>(x)};
648  }
649  inline
650  __device__ __host__
651  __half __float2half_rz(float x)
652  {
653  return __half_raw{static_cast<_Float16>(x)};
654  }
655  inline
656  __device__ __host__
657  __half __float2half_rd(float x)
658  {
659  return __half_raw{static_cast<_Float16>(x)};
660  }
661  inline
662  __device__ __host__
663  __half __float2half_ru(float x)
664  {
665  return __half_raw{static_cast<_Float16>(x)};
666  }
667  inline
668  __device__ __host__
669  __half2 __float2half2_rn(float x)
670  {
671  return __half2_raw{
672  _Float16_2{
673  static_cast<_Float16>(x), static_cast<_Float16>(x)}};
674  }
675  inline
676  __device__ __host__
677  __half2 __floats2half2_rn(float x, float y)
678  {
679  return __half2_raw{_Float16_2{
680  static_cast<_Float16>(x), static_cast<_Float16>(y)}};
681  }
682  inline
683  __device__ __host__
684  __half2 __float22half2_rn(float2 x)
685  {
686  return __floats2half2_rn(x.x, x.y);
687  }
688 
689  // half | half2 -> float
690  inline
691  __device__ __host__
692  float __half2float(__half x)
693  {
694  return static_cast<__half_raw>(x).data;
695  }
696  inline
697  __device__ __host__
698  float __low2float(__half2 x)
699  {
700  return static_cast<__half2_raw>(x).data.x;
701  }
702  inline
703  __device__ __host__
704  float __high2float(__half2 x)
705  {
706  return static_cast<__half2_raw>(x).data.y;
707  }
708  inline
709  __device__ __host__
710  float2 __half22float2(__half2 x)
711  {
712  return make_float2(
713  static_cast<__half2_raw>(x).data.x,
714  static_cast<__half2_raw>(x).data.y);
715  }
716 
717  // half -> int
718  inline
719  __device__
720  int __half2int_rn(__half x)
721  {
722  return static_cast<__half_raw>(x).data;
723  }
724  inline
725  __device__
726  int __half2int_rz(__half x)
727  {
728  return static_cast<__half_raw>(x).data;
729  }
730  inline
731  __device__
732  int __half2int_rd(__half x)
733  {
734  return static_cast<__half_raw>(x).data;
735  }
736  inline
737  __device__
738  int __half2int_ru(__half x)
739  {
740  return static_cast<__half_raw>(x).data;
741  }
742 
743  // int -> half
744  inline
745  __device__
746  __half __int2half_rn(int x)
747  {
748  return __half_raw{static_cast<_Float16>(x)};
749  }
750  inline
751  __device__
752  __half __int2half_rz(int x)
753  {
754  return __half_raw{static_cast<_Float16>(x)};
755  }
756  inline
757  __device__
758  __half __int2half_rd(int x)
759  {
760  return __half_raw{static_cast<_Float16>(x)};
761  }
762  inline
763  __device__
764  __half __int2half_ru(int x)
765  {
766  return __half_raw{static_cast<_Float16>(x)};
767  }
768 
769  // half -> short
770  inline
771  __device__
772  short __half2short_rn(__half x)
773  {
774  return static_cast<__half_raw>(x).data;
775  }
776  inline
777  __device__
778  short __half2short_rz(__half x)
779  {
780  return static_cast<__half_raw>(x).data;
781  }
782  inline
783  __device__
784  short __half2short_rd(__half x)
785  {
786  return static_cast<__half_raw>(x).data;
787  }
788  inline
789  __device__
790  short __half2short_ru(__half x)
791  {
792  return static_cast<__half_raw>(x).data;
793  }
794 
795  // short -> half
796  inline
797  __device__
798  __half __short2half_rn(short x)
799  {
800  return __half_raw{static_cast<_Float16>(x)};
801  }
802  inline
803  __device__
804  __half __short2half_rz(short x)
805  {
806  return __half_raw{static_cast<_Float16>(x)};
807  }
808  inline
809  __device__
810  __half __short2half_rd(short x)
811  {
812  return __half_raw{static_cast<_Float16>(x)};
813  }
814  inline
815  __device__
816  __half __short2half_ru(short x)
817  {
818  return __half_raw{static_cast<_Float16>(x)};
819  }
820 
821  // half -> long long
822  inline
823  __device__
824  long long __half2ll_rn(__half x)
825  {
826  return static_cast<__half_raw>(x).data;
827  }
828  inline
829  __device__
830  long long __half2ll_rz(__half x)
831  {
832  return static_cast<__half_raw>(x).data;
833  }
834  inline
835  __device__
836  long long __half2ll_rd(__half x)
837  {
838  return static_cast<__half_raw>(x).data;
839  }
840  inline
841  __device__
842  long long __half2ll_ru(__half x)
843  {
844  return static_cast<__half_raw>(x).data;
845  }
846 
847  // long long -> half
848  inline
849  __device__
850  __half __ll2half_rn(long long x)
851  {
852  return __half_raw{static_cast<_Float16>(x)};
853  }
854  inline
855  __device__
856  __half __ll2half_rz(long long x)
857  {
858  return __half_raw{static_cast<_Float16>(x)};
859  }
860  inline
861  __device__
862  __half __ll2half_rd(long long x)
863  {
864  return __half_raw{static_cast<_Float16>(x)};
865  }
866  inline
867  __device__
868  __half __ll2half_ru(long long x)
869  {
870  return __half_raw{static_cast<_Float16>(x)};
871  }
872 
873  // half -> unsigned int
874  inline
875  __device__
876  unsigned int __half2uint_rn(__half x)
877  {
878  return static_cast<__half_raw>(x).data;
879  }
880  inline
881  __device__
882  unsigned int __half2uint_rz(__half x)
883  {
884  return static_cast<__half_raw>(x).data;
885  }
886  inline
887  __device__
888  unsigned int __half2uint_rd(__half x)
889  {
890  return static_cast<__half_raw>(x).data;
891  }
892  inline
893  __device__
894  unsigned int __half2uint_ru(__half x)
895  {
896  return static_cast<__half_raw>(x).data;
897  }
898 
899  // unsigned int -> half
900  inline
901  __device__
902  __half __uint2half_rn(unsigned int x)
903  {
904  return __half_raw{static_cast<_Float16>(x)};
905  }
906  inline
907  __device__
908  __half __uint2half_rz(unsigned int x)
909  {
910  return __half_raw{static_cast<_Float16>(x)};
911  }
912  inline
913  __device__
914  __half __uint2half_rd(unsigned int x)
915  {
916  return __half_raw{static_cast<_Float16>(x)};
917  }
918  inline
919  __device__
920  __half __uint2half_ru(unsigned int x)
921  {
922  return __half_raw{static_cast<_Float16>(x)};
923  }
924 
925  // half -> unsigned short
926  inline
927  __device__
928  unsigned short __half2ushort_rn(__half x)
929  {
930  return static_cast<__half_raw>(x).data;
931  }
932  inline
933  __device__
934  unsigned short __half2ushort_rz(__half x)
935  {
936  return static_cast<__half_raw>(x).data;
937  }
938  inline
939  __device__
940  unsigned short __half2ushort_rd(__half x)
941  {
942  return static_cast<__half_raw>(x).data;
943  }
944  inline
945  __device__
946  unsigned short __half2ushort_ru(__half x)
947  {
948  return static_cast<__half_raw>(x).data;
949  }
950 
951  // unsigned short -> half
952  inline
953  __device__
954  __half __ushort2half_rn(unsigned short x)
955  {
956  return __half_raw{static_cast<_Float16>(x)};
957  }
958  inline
959  __device__
960  __half __ushort2half_rz(unsigned short x)
961  {
962  return __half_raw{static_cast<_Float16>(x)};
963  }
964  inline
965  __device__
966  __half __ushort2half_rd(unsigned short x)
967  {
968  return __half_raw{static_cast<_Float16>(x)};
969  }
970  inline
971  __device__
972  __half __ushort2half_ru(unsigned short x)
973  {
974  return __half_raw{static_cast<_Float16>(x)};
975  }
976 
977  // half -> unsigned long long
978  inline
979  __device__
980  unsigned long long __half2ull_rn(__half x)
981  {
982  return static_cast<__half_raw>(x).data;
983  }
984  inline
985  __device__
986  unsigned long long __half2ull_rz(__half x)
987  {
988  return static_cast<__half_raw>(x).data;
989  }
990  inline
991  __device__
992  unsigned long long __half2ull_rd(__half x)
993  {
994  return static_cast<__half_raw>(x).data;
995  }
996  inline
997  __device__
998  unsigned long long __half2ull_ru(__half x)
999  {
1000  return static_cast<__half_raw>(x).data;
1001  }
1002 
1003  // unsigned long long -> half
1004  inline
1005  __device__
1006  __half __ull2half_rn(unsigned long long x)
1007  {
1008  return __half_raw{static_cast<_Float16>(x)};
1009  }
1010  inline
1011  __device__
1012  __half __ull2half_rz(unsigned long long x)
1013  {
1014  return __half_raw{static_cast<_Float16>(x)};
1015  }
1016  inline
1017  __device__
1018  __half __ull2half_rd(unsigned long long x)
1019  {
1020  return __half_raw{static_cast<_Float16>(x)};
1021  }
1022  inline
1023  __device__
1024  __half __ull2half_ru(unsigned long long x)
1025  {
1026  return __half_raw{static_cast<_Float16>(x)};
1027  }
1028 
1029  // Load primitives
1030  inline
1031  __device__
1032  __half __ldg(const __half* ptr) { return *ptr; }
1033  inline
1034  __device__
1035  __half __ldcg(const __half* ptr) { return *ptr; }
1036  inline
1037  __device__
1038  __half __ldca(const __half* ptr) { return *ptr; }
1039  inline
1040  __device__
1041  __half __ldcs(const __half* ptr) { return *ptr; }
1042 
1043  inline
1044  __device__
1045  __half2 __ldg(const __half2* ptr) { return *ptr; }
1046  inline
1047  __device__
1048  __half2 __ldcg(const __half2* ptr) { return *ptr; }
1049  inline
1050  __device__
1051  __half2 __ldca(const __half2* ptr) { return *ptr; }
1052  inline
1053  __device__
1054  __half2 __ldcs(const __half2* ptr) { return *ptr; }
1055 
1056  // Relations
1057  inline
1058  __device__
1059  bool __heq(__half x, __half y)
1060  {
1061  return static_cast<__half_raw>(x).data ==
1062  static_cast<__half_raw>(y).data;
1063  }
1064  inline
1065  __device__
1066  bool __hne(__half x, __half y)
1067  {
1068  return static_cast<__half_raw>(x).data !=
1069  static_cast<__half_raw>(y).data;
1070  }
1071  inline
1072  __device__
1073  bool __hle(__half x, __half y)
1074  {
1075  return static_cast<__half_raw>(x).data <=
1076  static_cast<__half_raw>(y).data;
1077  }
1078  inline
1079  __device__
1080  bool __hge(__half x, __half y)
1081  {
1082  return static_cast<__half_raw>(x).data >=
1083  static_cast<__half_raw>(y).data;
1084  }
1085  inline
1086  __device__
1087  bool __hlt(__half x, __half y)
1088  {
1089  return static_cast<__half_raw>(x).data <
1090  static_cast<__half_raw>(y).data;
1091  }
1092  inline
1093  __device__
1094  bool __hgt(__half x, __half y)
1095  {
1096  return static_cast<__half_raw>(x).data >
1097  static_cast<__half_raw>(y).data;
1098  }
1099  inline
1100  __device__
1101  bool __hequ(__half x, __half y) { return __heq(x, y); }
1102  inline
1103  __device__
1104  bool __hneu(__half x, __half y) { return __hne(x, y); }
1105  inline
1106  __device__
1107  bool __hleu(__half x, __half y) { return __hle(x, y); }
1108  inline
1109  __device__
1110  bool __hgeu(__half x, __half y) { return __hge(x, y); }
1111  inline
1112  __device__
1113  bool __hltu(__half x, __half y) { return __hlt(x, y); }
1114  inline
1115  __device__
1116  bool __hgtu(__half x, __half y) { return __hgt(x, y); }
1117 
1118  inline
1119  __device__
1120  __half2 __heq2(__half2 x, __half2 y)
1121  {
1122  auto r = static_cast<__half2_raw>(x).data ==
1123  static_cast<__half2_raw>(y).data;
1124  return __half2_raw{_Float16_2{
1125  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1126  }
1127  inline
1128  __device__
1129  __half2 __hne2(__half2 x, __half2 y)
1130  {
1131  auto r = static_cast<__half2_raw>(x).data !=
1132  static_cast<__half2_raw>(y).data;
1133  return __half2_raw{_Float16_2{
1134  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1135  }
1136  inline
1137  __device__
1138  __half2 __hle2(__half2 x, __half2 y)
1139  {
1140  auto r = static_cast<__half2_raw>(x).data <=
1141  static_cast<__half2_raw>(y).data;
1142  return __half2_raw{_Float16_2{
1143  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1144  }
1145  inline
1146  __device__
1147  __half2 __hge2(__half2 x, __half2 y)
1148  {
1149  auto r = static_cast<__half2_raw>(x).data >=
1150  static_cast<__half2_raw>(y).data;
1151  return __half2_raw{_Float16_2{
1152  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1153  }
1154  inline
1155  __device__
1156  __half2 __hlt2(__half2 x, __half2 y)
1157  {
1158  auto r = static_cast<__half2_raw>(x).data <
1159  static_cast<__half2_raw>(y).data;
1160  return __half2_raw{_Float16_2{
1161  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1162  }
1163  inline
1164  __device__
1165  __half2 __hgt2(__half2 x, __half2 y)
1166  {
1167  auto r = static_cast<__half2_raw>(x).data >
1168  static_cast<__half2_raw>(y).data;
1169  return __half2_raw{_Float16_2{
1170  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1171  }
1172  inline
1173  __device__
1174  __half2 __hequ2(__half2 x, __half2 y) { return __heq2(x, y); }
1175  inline
1176  __device__
1177  __half2 __hneu2(__half2 x, __half2 y) { return __hne2(x, y); }
1178  inline
1179  __device__
1180  __half2 __hleu2(__half2 x, __half2 y) { return __hle2(x, y); }
1181  inline
1182  __device__
1183  __half2 __hgeu2(__half2 x, __half2 y) { return __hge2(x, y); }
1184  inline
1185  __device__
1186  __half2 __hltu2(__half2 x, __half2 y) { return __hlt2(x, y); }
1187  inline
1188  __device__
1189  __half2 __hgtu2(__half2 x, __half2 y) { return __hgt2(x, y); }
1190 
1191  inline
1192  __device__
1193  bool __hbeq2(__half2 x, __half2 y)
1194  {
1195  auto r = static_cast<__half2_raw>(__heq2(x, y));
1196  return r.data.x != 0 && r.data.y != 0;
1197  }
1198  inline
1199  __device__
1200  bool __hbne2(__half2 x, __half2 y)
1201  {
1202  auto r = static_cast<__half2_raw>(__hne2(x, y));
1203  return r.data.x != 0 && r.data.y != 0;
1204  }
1205  inline
1206  __device__
1207  bool __hble2(__half2 x, __half2 y)
1208  {
1209  auto r = static_cast<__half2_raw>(__hle2(x, y));
1210  return r.data.x != 0 && r.data.y != 0;
1211  }
1212  inline
1213  __device__
1214  bool __hbge2(__half2 x, __half2 y)
1215  {
1216  auto r = static_cast<__half2_raw>(__hge2(x, y));
1217  return r.data.x != 0 && r.data.y != 0;
1218  }
1219  inline
1220  __device__
1221  bool __hblt2(__half2 x, __half2 y)
1222  {
1223  auto r = static_cast<__half2_raw>(__hlt2(x, y));
1224  return r.data.x != 0 && r.data.y != 0;
1225  }
1226  inline
1227  __device__
1228  bool __hbgt2(__half2 x, __half2 y)
1229  {
1230  auto r = static_cast<__half2_raw>(__hgt2(x, y));
1231  return r.data.x != 0 && r.data.y != 0;
1232  }
1233  inline
1234  __device__
1235  bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); }
1236  inline
1237  __device__
1238  bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); }
1239  inline
1240  __device__
1241  bool __hbleu2(__half2 x, __half2 y) { return __hble2(x, y); }
1242  inline
1243  __device__
1244  bool __hbgeu2(__half2 x, __half2 y) { return __hbge2(x, y); }
1245  inline
1246  __device__
1247  bool __hbltu2(__half2 x, __half2 y) { return __hblt2(x, y); }
1248  inline
1249  __device__
1250  bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); }
1251 
1252  // Arithmetic
1253  inline
1254  __device__
1255  __half __clamp_01(__half x)
1256  {
1257  auto r = static_cast<__half_raw>(x);
1258 
1259  if (__hlt(x, __half_raw{0})) return __half_raw{0};
1260  if (__hlt(__half_raw{1}, x)) return __half_raw{1};
1261  return r;
1262  }
1263 
1264  inline
1265  __device__
1266  __half __hadd(__half x, __half y)
1267  {
1268  return __half_raw{
1269  static_cast<__half_raw>(x).data +
1270  static_cast<__half_raw>(y).data};
1271  }
1272  inline
1273  __device__
1274  __half __hsub(__half x, __half y)
1275  {
1276  return __half_raw{
1277  static_cast<__half_raw>(x).data -
1278  static_cast<__half_raw>(y).data};
1279  }
1280  inline
1281  __device__
1282  __half __hmul(__half x, __half y)
1283  {
1284  return __half_raw{
1285  static_cast<__half_raw>(x).data *
1286  static_cast<__half_raw>(y).data};
1287  }
1288  inline
1289  __device__
1290  __half __hadd_sat(__half x, __half y)
1291  {
1292  return __clamp_01(__hadd(x, y));
1293  }
1294  inline
1295  __device__
1296  __half __hsub_sat(__half x, __half y)
1297  {
1298  return __clamp_01(__hsub(x, y));
1299  }
1300  inline
1301  __device__
1302  __half __hmul_sat(__half x, __half y)
1303  {
1304  return __clamp_01(__hmul(x, y));
1305  }
1306  inline
1307  __device__
1308  __half __hfma(__half x, __half y, __half z)
1309  {
1310  return __half_raw{__ocml_fma_f16(
1311  static_cast<__half_raw>(x).data,
1312  static_cast<__half_raw>(y).data,
1313  static_cast<__half_raw>(z).data)};
1314  }
1315  inline
1316  __device__
1317  __half __hfma_sat(__half x, __half y, __half z)
1318  {
1319  return __clamp_01(__hfma(x, y, z));
1320  }
1321  inline
1322  __device__
1323  __half __hdiv(__half x, __half y)
1324  {
1325  return __half_raw{
1326  static_cast<__half_raw>(x).data /
1327  static_cast<__half_raw>(y).data};
1328  }
1329 
1330  inline
1331  __device__
1332  __half2 __hadd2(__half2 x, __half2 y)
1333  {
1334  return __half2_raw{
1335  static_cast<__half2_raw>(x).data +
1336  static_cast<__half2_raw>(y).data};
1337  }
1338  inline
1339  __device__
1340  __half2 __hsub2(__half2 x, __half2 y)
1341  {
1342  return __half2_raw{
1343  static_cast<__half2_raw>(x).data -
1344  static_cast<__half2_raw>(y).data};
1345  }
1346  inline
1347  __device__
1348  __half2 __hmul2(__half2 x, __half2 y)
1349  {
1350  return __half2_raw{
1351  static_cast<__half2_raw>(x).data *
1352  static_cast<__half2_raw>(y).data};
1353  }
1354  inline
1355  __device__
1356  __half2 __hadd2_sat(__half2 x, __half2 y)
1357  {
1358  auto r = static_cast<__half2_raw>(__hadd2(x, y));
1359  return __half2{
1360  __clamp_01(__half_raw{r.data.x}),
1361  __clamp_01(__half_raw{r.data.y})};
1362  }
1363  inline
1364  __device__
1365  __half2 __hsub2_sat(__half2 x, __half2 y)
1366  {
1367  auto r = static_cast<__half2_raw>(__hsub2(x, y));
1368  return __half2{
1369  __clamp_01(__half_raw{r.data.x}),
1370  __clamp_01(__half_raw{r.data.y})};
1371  }
1372  inline
1373  __device__
1374  __half2 __hmul2_sat(__half2 x, __half2 y)
1375  {
1376  auto r = static_cast<__half2_raw>(__hmul2(x, y));
1377  return __half2{
1378  __clamp_01(__half_raw{r.data.x}),
1379  __clamp_01(__half_raw{r.data.y})};
1380  }
1381  inline
1382  __device__
1383  __half2 __hfma2(__half2 x, __half2 y, __half2 z)
1384  {
1385  return __half2_raw{__ocml_fma_2f16(x, y, z)};
1386  }
1387  inline
1388  __device__
1389  __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z)
1390  {
1391  auto r = static_cast<__half2_raw>(__hfma2(x, y, z));
1392  return __half2{
1393  __clamp_01(__half_raw{r.data.x}),
1394  __clamp_01(__half_raw{r.data.y})};
1395  }
1396  inline
1397  __device__
1398  __half2 __h2div(__half2 x, __half2 y)
1399  {
1400  return __half2_raw{
1401  static_cast<__half2_raw>(x).data /
1402  static_cast<__half2_raw>(y).data};
1403  }
1404 
1405  // Math functions
1406  #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
1407  inline
1408  __device__
1409  float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) {
1410  return __ockl_fdot2(static_cast<__half2_raw>(a).data,
1411  static_cast<__half2_raw>(b).data,
1412  c, saturate);
1413  }
1414  #endif
1415  inline
1416  __device__
1417  __half htrunc(__half x)
1418  {
1419  return __half_raw{
1420  __ocml_trunc_f16(static_cast<__half_raw>(x).data)};
1421  }
1422  inline
1423  __device__
1424  __half hceil(__half x)
1425  {
1426  return __half_raw{
1427  __ocml_ceil_f16(static_cast<__half_raw>(x).data)};
1428  }
1429  inline
1430  __device__
1431  __half hfloor(__half x)
1432  {
1433  return __half_raw{
1434  __ocml_floor_f16(static_cast<__half_raw>(x).data)};
1435  }
1436  inline
1437  __device__
1438  __half hrint(__half x)
1439  {
1440  return __half_raw{
1441  __ocml_rint_f16(static_cast<__half_raw>(x).data)};
1442  }
1443  inline
1444  __device__
1445  __half hsin(__half x)
1446  {
1447  return __half_raw{
1448  __ocml_sin_f16(static_cast<__half_raw>(x).data)};
1449  }
1450  inline
1451  __device__
1452  __half hcos(__half x)
1453  {
1454  return __half_raw{
1455  __ocml_cos_f16(static_cast<__half_raw>(x).data)};
1456  }
1457  inline
1458  __device__
1459  __half hexp(__half x)
1460  {
1461  return __half_raw{
1462  __ocml_exp_f16(static_cast<__half_raw>(x).data)};
1463  }
1464  inline
1465  __device__
1466  __half hexp2(__half x)
1467  {
1468  return __half_raw{
1469  __ocml_exp2_f16(static_cast<__half_raw>(x).data)};
1470  }
1471  inline
1472  __device__
1473  __half hexp10(__half x)
1474  {
1475  return __half_raw{
1476  __ocml_exp10_f16(static_cast<__half_raw>(x).data)};
1477  }
1478  inline
1479  __device__
1480  __half hlog2(__half x)
1481  {
1482  return __half_raw{
1483  __ocml_log2_f16(static_cast<__half_raw>(x).data)};
1484  }
1485  inline
1486  __device__
1487  __half hlog(__half x)
1488  {
1489  return __half_raw{
1490  __ocml_log_f16(static_cast<__half_raw>(x).data)};
1491  }
1492  inline
1493  __device__
1494  __half hlog10(__half x)
1495  {
1496  return __half_raw{
1497  __ocml_log10_f16(static_cast<__half_raw>(x).data)};
1498  }
1499  inline
1500  __device__
1501  __half hrcp(__half x)
1502  {
1503  return __half_raw{
1504  __llvm_amdgcn_rcp_f16(static_cast<__half_raw>(x).data)};
1505  }
1506  inline
1507  __device__
1508  __half hrsqrt(__half x)
1509  {
1510  return __half_raw{
1511  __ocml_rsqrt_f16(static_cast<__half_raw>(x).data)};
1512  }
1513  inline
1514  __device__
1515  __half hsqrt(__half x)
1516  {
1517  return __half_raw{
1518  __ocml_sqrt_f16(static_cast<__half_raw>(x).data)};
1519  }
1520  inline
1521  __device__
1522  bool __hisinf(__half x)
1523  {
1524  return __ocml_isinf_f16(static_cast<__half_raw>(x).data);
1525  }
1526  inline
1527  __device__
1528  bool __hisnan(__half x)
1529  {
1530  return __ocml_isnan_f16(static_cast<__half_raw>(x).data);
1531  }
1532  inline
1533  __device__
1534  __half __hneg(__half x)
1535  {
1536  return __half_raw{-static_cast<__half_raw>(x).data};
1537  }
1538 
1539  inline
1540  __device__
1541  __half2 h2trunc(__half2 x)
1542  {
1543  return __half2_raw{__ocml_trunc_2f16(x)};
1544  }
1545  inline
1546  __device__
1547  __half2 h2ceil(__half2 x)
1548  {
1549  return __half2_raw{__ocml_ceil_2f16(x)};
1550  }
1551  inline
1552  __device__
1553  __half2 h2floor(__half2 x)
1554  {
1555  return __half2_raw{__ocml_floor_2f16(x)};
1556  }
1557  inline
1558  __device__
1559  __half2 h2rint(__half2 x)
1560  {
1561  return __half2_raw{__ocml_rint_2f16(x)};
1562  }
1563  inline
1564  __device__
1565  __half2 h2sin(__half2 x)
1566  {
1567  return __half2_raw{__ocml_sin_2f16(x)};
1568  }
1569  inline
1570  __device__
1571  __half2 h2cos(__half2 x)
1572  {
1573  return __half2_raw{__ocml_cos_2f16(x)};
1574  }
1575  inline
1576  __device__
1577  __half2 h2exp(__half2 x)
1578  {
1579  return __half2_raw{__ocml_exp_2f16(x)};
1580  }
1581  inline
1582  __device__
1583  __half2 h2exp2(__half2 x)
1584  {
1585  return __half2_raw{__ocml_exp2_2f16(x)};
1586  }
1587  inline
1588  __device__
1589  __half2 h2exp10(__half2 x)
1590  {
1591  return __half2_raw{__ocml_exp10_2f16(x)};
1592  }
1593  inline
1594  __device__
1595  __half2 h2log2(__half2 x)
1596  {
1597  return __half2_raw{__ocml_log2_2f16(x)};
1598  }
1599  inline
1600  __device__
1601  __half2 h2log(__half2 x) { return __ocml_log_2f16(x); }
1602  inline
1603  __device__
1604  __half2 h2log10(__half2 x) { return __ocml_log10_2f16(x); }
1605  inline
1606  __device__
1607  __half2 h2rcp(__half2 x) { return __llvm_amdgcn_rcp_2f16(x); }
1608  inline
1609  __device__
1610  __half2 h2rsqrt(__half2 x) { return __ocml_rsqrt_2f16(x); }
1611  inline
1612  __device__
1613  __half2 h2sqrt(__half2 x) { return __ocml_sqrt_2f16(x); }
1614  inline
1615  __device__
1616  __half2 __hisinf2(__half2 x)
1617  {
1618  auto r = __ocml_isinf_2f16(x);
1619  return __half2_raw{_Float16_2{
1620  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1621  }
1622  inline
1623  __device__
1624  __half2 __hisnan2(__half2 x)
1625  {
1626  auto r = __ocml_isnan_2f16(x);
1627  return __half2_raw{_Float16_2{
1628  static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
1629  }
1630  inline
1631  __device__
1632  __half2 __hneg2(__half2 x)
1633  {
1634  return __half2_raw{-static_cast<__half2_raw>(x).data};
1635  }
1636  } // Anonymous namespace.
1637 
1638  #if !defined(HIP_NO_HALF)
1639  using half = __half;
1640  using half2 = __half2;
1641  #endif
1642  #endif // defined(__cplusplus)
1643 #elif defined(__GNUC__)
1644  #include "hip_fp16_gcc.h"
1645 #endif // !defined(__clang__) && defined(__GNUC__)
TODO-doc.
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Copies the memory address of symbol symbolName to devPtr.
Definition: hip_fp16_math_fwd.h:53
Definition: hip_fp16_gcc.h:11
#define __host__
Definition: host_defines.h:41
Definition: hip_fp16_gcc.h:7