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