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