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