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