30#ifndef _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP8_H_
31#define _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP8_H_
33#if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) && __HIP_DEVICE_COMPILE__
34#define HIP_FP8_CVT_FAST_PATH 1
36#define HIP_FP8_CVT_FAST_PATH 0
39#if !defined(__HIPCC_RTC__)
40#include <hip/amd_detail/amd_hip_common.h>
44#include "amd_hip_vector_types.h"
45#include "amd_hip_fp16.h"
50#if defined(__HIPCC_RTC__)
51#define __FP8_HOST_DEVICE__ __device__
52#define __FP8_HOST_DEVICE_STATIC__ __FP8_HOST_DEVICE__ static
54#define __FP8_HOST_DEVICE__ __host__ __device__
55#define __FP8_HOST_DEVICE_STATIC__ __FP8_HOST_DEVICE__ static inline
58#if !defined(__HIPCC_RTC__)
59static_assert(CHAR_BIT == 8,
"byte size should be of 8 bits");
61static_assert(
sizeof(
unsigned char) == 1);
62static_assert(
sizeof(
unsigned short int) == 2);
63static_assert(
sizeof(
unsigned int) == 4);
105template <
typename T,
bool negative_zero_nan>
106__FP8_HOST_DEVICE_STATIC__
__hip_fp8_storage_t cast_to_f8(T _x,
int wm,
int we,
bool clip =
false,
108 unsigned int rng = 0) {
109 constexpr bool is_half = __hip_internal::is_same<T, _Float16>::value;
110 constexpr bool is_float = __hip_internal::is_same<T, float>::value;
111 constexpr bool is_double = __hip_internal::is_same<T, double>::value;
112 static_assert(is_half || is_float || is_double,
"Only half, float and double can be cast to f8");
114 const int mfmt = (
sizeof(T) == 8) ? 52 : ((
sizeof(T) == 4) ? 23 : 10);
115 unsigned long long x;
118 x =
reinterpret_cast<unsigned long long&
>(_x);
119 else if (
sizeof(T) == 4)
120 x =
reinterpret_cast<unsigned int&
>(_x);
122 x =
reinterpret_cast<unsigned short int&
>(_x);
125 unsigned long long head, mantissa;
129 if (
sizeof(T) == 8) {
130 head = x & 0xFFF0000000000000ull;
131 mantissa = x & 0xFFFFFFFFFFFFFull;
132 exponent = (head >> 52) & 0x7FF;
135 }
else if (
sizeof(T) == 4) {
136 head = x & 0xFF800000;
137 mantissa = x & 0x7FFFFF;
138 exponent = (head >> 23) & 0xFF;
143 mantissa = x & 0x3FF;
144 exponent = (head >> 10) & 0x1F;
149 unsigned int signed_inf = (sign << 7) + (((1 << we) - 1) << wm);
152 if (negative_zero_nan) {
153 if (
sizeof(T) == 8) {
154 if ((x & 0x7FF0000000000000ull) == 0x7FF0000000000000ull)
return 0x80;
155 }
else if (
sizeof(T) == 4) {
156 if ((x & 0x7F800000) == 0x7F800000)
return 0x80;
158 if ((x & 0x7C00) == 0x7C00)
return 0x80;
161 if (
sizeof(T) == 8) {
162 if ((x & 0x7FF0000000000000ull) == 0x7FF0000000000000ull)
163 return signed_inf + (mantissa != 0 ? 1 : 0);
164 }
else if (
sizeof(T) == 4) {
165 if ((x & 0x7F800000) == 0x7F800000)
return signed_inf + (mantissa != 0 ? 1 : 0);
167 if ((x & 0x7C00) == 0x7C00)
return signed_inf + (mantissa != 0 ? 1 : 0);
182 const int f8_bias = (1 << (we - 1)) - 1 + (negative_zero_nan ? 1 : 0);
183 const int f8_denormal_act_exponent = 1 - f8_bias;
188 int act_exponent, f8_exponent, exponent_diff;
197 act_exponent = exponent - bias + 1;
198 exponent_diff = f8_denormal_act_exponent -
201 act_exponent = exponent - bias;
202 if (act_exponent <= f8_denormal_act_exponent) {
208 exponent_diff = f8_denormal_act_exponent - act_exponent;
213 mantissa += (1ull << mfmt);
216 bool midpoint = (mantissa & ((1ull << (mfmt - wm + exponent_diff)) - 1)) ==
217 (1ull << (mfmt - wm + exponent_diff - 1));
224 if (exponent_diff > 0)
225 mantissa >>= exponent_diff;
226 else if (exponent_diff == -1)
227 mantissa <<= -exponent_diff;
228 bool implicit_one = mantissa & (1ull << mfmt);
231 (act_exponent + exponent_diff) + f8_bias - (implicit_one ? 0 : 1);
234 unsigned long long drop_mask = (1ull << (mfmt - wm)) - 1;
236 mantissa & (1ull << (mfmt - wm));
238 (stoch ? rng : (midpoint ? (odd ? mantissa : mantissa - 1ull) : mantissa)) & drop_mask;
241 if (f8_exponent == 0) {
242 if ((1ull << mfmt) & mantissa) {
246 if ((1ull << (mfmt + 1)) & mantissa) {
252 mantissa >>= (mfmt - wm);
255 const int max_exp = (1 << we) - (negative_zero_nan ? 1 : 2);
256 if (f8_exponent > max_exp) {
258 mantissa = (1 << wm) - 1;
259 f8_exponent = max_exp;
265 if (f8_exponent == 0 && mantissa == 0)
return negative_zero_nan ? 0 : (sign << 7);
266 mantissa &= (1 << wm) - 1;
267 return (sign << 7) | (f8_exponent << wm) | mantissa;
273template <
typename T,
bool negative_zero_nan>
275 constexpr bool is_half = __hip_internal::is_same<T, _Float16>::value;
276 constexpr bool is_float = __hip_internal::is_same<T, float>::value;
277 constexpr bool is_double = __hip_internal::is_same<T, double>::value;
278 static_assert(is_half || is_float || is_double,
"only half, float and double are supported");
280 constexpr int weo = is_half ? 5 : (is_float ? 8 : 11);
281 constexpr int wmo = is_half ? 10 : (is_float ? 23 : 52);
283 T fInf, fNegInf, fNaN, fNeg0;
285 const unsigned short int ihInf = 0x7C00;
286 const unsigned short int ihNegInf = 0xFC00;
287 const unsigned short int ihNaN = 0x7C01;
288 const unsigned short int ihNeg0 = 0x8000;
289 fInf =
reinterpret_cast<const _Float16&
>(ihInf);
290 fNegInf =
reinterpret_cast<const _Float16&
>(ihNegInf);
291 fNaN =
reinterpret_cast<const _Float16&
>(ihNaN);
292 fNeg0 =
reinterpret_cast<const _Float16&
>(ihNeg0);
293 }
else if (is_float) {
294 const unsigned int ifInf = 0x7F800000;
295 const unsigned int ifNegInf = 0xFF800000;
296 const unsigned int ifNaN = 0x7F800001;
297 const unsigned int ifNeg0 = 0x80000000;
298 fInf =
reinterpret_cast<const float&
>(ifInf);
299 fNegInf =
reinterpret_cast<const float&
>(ifNegInf);
300 fNaN =
reinterpret_cast<const float&
>(ifNaN);
301 fNeg0 =
reinterpret_cast<const float&
>(ifNeg0);
302 }
else if (is_double) {
303 const unsigned long long ifInf = 0x7FF0000000000000ull;
304 const unsigned long long ifNegInf = 0xFFF0000000000000ull;
305 const unsigned long long ifNaN = 0x7FF0000000000001ull;
306 const unsigned long long ifNeg0 = 0x8000000000000000ull;
307 fInf =
reinterpret_cast<const double&
>(ifInf);
308 fNegInf =
reinterpret_cast<const double&
>(ifNegInf);
309 fNaN =
reinterpret_cast<const double&
>(ifNaN);
310 fNeg0 =
reinterpret_cast<const double&
>(ifNeg0);
317 unsigned long long sign = x >> 7;
318 unsigned long long mantissa = x & ((1 << wm) - 1);
319 int exponent = (x & 0x7F) >> wm;
320 if (negative_zero_nan) {
321 if (x == 0x80)
return fNaN;
323 if (x == 0x80)
return fNeg0;
324 if (exponent == ((1 << we) - 1))
return (mantissa == 0) ? (sign ? fNegInf : fInf) : fNaN;
327 typename __hip_internal::conditional<
328 sizeof(T) == 2,
unsigned short int,
329 typename __hip_internal::conditional<
sizeof(T) == 4,
unsigned int,
330 unsigned long long>::type>::type retval;
332 if (we == 5 && is_half && !negative_zero_nan) {
334 return reinterpret_cast<const T&
>(retval);
337 const int exp_low_cutoff = (1 << (weo - 1)) - (1 << (we - 1)) + 1 - (negative_zero_nan ? 1 : 0);
341#if __HIP_DEVICE_COMPILE__
343 int sh = 1 + __clz(mantissa) - (32 - wm);
345 int sh = 1 + __builtin_clz(mantissa) - (32 - wm);
349 mantissa &= ((1ull << wm) - 1);
351 exponent += exp_low_cutoff - 1;
352 mantissa <<= wmo - wm;
356 mantissa |= 1 << wmo;
357 mantissa >>= 1 - exponent;
362 retval = (sign << 15) | (exponent << 10) | mantissa;
363 else if (
sizeof(T) == 4)
364 retval = (sign << 31) | (exponent << 23) | mantissa;
366 retval = (sign << 63) | (static_cast<unsigned long long>(exponent) << 52) | mantissa;
367 return reinterpret_cast<const T&
>(retval);
370#if HIP_FP8_CVT_FAST_PATH
373template <
bool stochastic_rounding = false>
376 unsigned int rng = 0) {
381 unsigned char i8val[4];
384 unsigned int ival = 0;
389 if ((val.i32val & 0x7F800000) != 0x7F800000) {
390 val.fval = __builtin_amdgcn_fmed3f(val.fval, 240.0, -240.0);
393 if ((val.i32val & 0x7F800000) != 0x7F800000) {
394 val.fval = __builtin_amdgcn_fmed3f(val.fval, 57344.0, -57344.0);
399 if (stochastic_rounding) {
401 ? __builtin_amdgcn_cvt_sr_fp8_f32(val.fval, rng, ival, 0)
402 : __builtin_amdgcn_cvt_sr_bf8_f32(val.fval, rng, ival, 0);
404 i8data = val.i8val[0];
407 ? __builtin_amdgcn_cvt_pk_fp8_f32(val.fval, val.fval, ival,
false)
408 : __builtin_amdgcn_cvt_pk_bf8_f32(val.fval, val.fval, ival,
false);
410 i8data = val.i8val[0];
418 static_assert(
sizeof(
float2) ==
sizeof(
unsigned int[2]));
419 static_assert(
sizeof(
float2) ==
sizeof(
unsigned short[4]));
421 unsigned int i32val[2];
422 unsigned short i16val[4];
428 if ((f2val.i32val[0] & 0x7F800000) != 0x7F800000) {
429 f2val.fval.x = __builtin_amdgcn_fmed3f(f2val.fval.x, 240.0, -240.0);
431 if ((f2val.i32val[1] & 0x7F800000) != 0x7F800000) {
432 f2val.fval.y = __builtin_amdgcn_fmed3f(f2val.fval.x, 240.0, -240.0);
437 ? __builtin_amdgcn_cvt_pk_fp8_f32(v.x, v.y, 0,
false)
438 : __builtin_amdgcn_cvt_pk_bf8_f32(v.x, v.y, 0,
false);
447 unsigned char i8val[4];
451 float fval = interpret ==
__HIP_E4M3_FNUZ ? __builtin_amdgcn_cvt_f32_fp8(val.i32val, 0)
452 : __builtin_amdgcn_cvt_f32_bf8(val.i32val, 0);
460 unsigned short i16val[2];
464 auto f2 = interpret ==
__HIP_E4M3_FNUZ ? __builtin_amdgcn_cvt_pk_f32_fp8(val.i32val,
false)
465 : __builtin_amdgcn_cvt_pk_f32_bf8(val.i32val,
false);
466 return float2{f2[0], f2[1]};
474 return static_cast<unsigned char>(a) == 0x80;
488#if HIP_FP8_CVT_FAST_PATH
489 return internal::cast_to_f8_from_f32<false>(f, sat ==
__HIP_SATFINITE, type);
493 return internal::cast_to_f8<float, true>(f, wm, we, sat ==
__HIP_SATFINITE);
507#if HIP_FP8_CVT_FAST_PATH
508 return internal::cast_to_f8x2_from_f32x2(f2, sat ==
__HIP_SATFINITE, type);
528 return internal::cast_to_f8<double, true>(d, wm, we, sat ==
__HIP_SATFINITE);
557 float fval = __hip_bfloat16(hr);
572 float2 f2 = __hip_bfloat162(hr);
587 return __half_raw{internal::cast_from_f8<_Float16, true>(x, wm, we)};
599 __half2 ret(
static_cast<__half
>(
640 constexpr static unsigned int __we = 4;
641 constexpr static unsigned int __wm = 3;
649 __default_interpret)) {}
654 __default_interpret)) {}
659 __default_interpret)) {}
664 __default_interpret)) {}
669 __default_interpret)) {}
674 __default_interpret)) {}
687 __default_interpret)) {}
692 __default_interpret)) {}
698 __FP8_HOST_DEVICE__
operator __half()
const {
703 __FP8_HOST_DEVICE__
operator __hip_bfloat16()
const {
705 return __hip_bfloat16(f);
709 __FP8_HOST_DEVICE__
operator bool()
const {
711 return !(
static_cast<unsigned short>(__x) == 0);
715 __FP8_HOST_DEVICE__
operator char()
const {
716 if (internal::hip_fp8_fnuz_is_nan(__x)) {
720 auto fval = internal::cast_from_f8<float, true>(__x, __wm, __we);
721 auto llval =
static_cast<long long>(fval);
722 if (llval <= CHAR_MIN) {
724 }
else if (llval >= CHAR_MAX) {
727 return static_cast<char>(fval);
731 __FP8_HOST_DEVICE__
operator double()
const {
732 return internal::cast_from_f8<double, true>(__x, __wm, __we);
736 __FP8_HOST_DEVICE__
operator float()
const {
737#if HIP_FP8_CVT_FAST_PATH
738 return internal::cast_to_f32_from_f8(__x, __default_interpret);
740 return internal::cast_from_f8<float, true>(__x, __wm, __we);
745 __FP8_HOST_DEVICE__
operator int()
const {
746 if (internal::hip_fp8_fnuz_is_nan(__x)) {
751 return static_cast<int>(fval);
755 __FP8_HOST_DEVICE__
operator long int()
const {
756 if (internal::hip_fp8_fnuz_is_nan(__x)) {
761 return static_cast<long>(fval);
765 __FP8_HOST_DEVICE__
operator long long int()
const {
766 if (internal::hip_fp8_fnuz_is_nan(__x)) {
771 return static_cast<long long>(fval);
775 __FP8_HOST_DEVICE__
operator short int()
const {
776 if (internal::hip_fp8_fnuz_is_nan(__x)) {
781 auto llval =
static_cast<long long>(fval);
782 if (llval <= SHRT_MIN) {
784 }
else if (llval >= SHRT_MAX) {
787 return static_cast<short>(fval);
791 __FP8_HOST_DEVICE__
operator signed char()
const {
792 if (internal::hip_fp8_fnuz_is_nan(__x)) {
797 auto llval =
static_cast<long long>(fval);
798 if (llval <= SCHAR_MIN) {
800 }
else if (llval >= SCHAR_MAX) {
803 return static_cast<signed char>(fval);
807 __FP8_HOST_DEVICE__
operator unsigned char()
const {
808 if (internal::hip_fp8_fnuz_is_nan(__x)) {
813 auto llval =
static_cast<long long>(fval);
816 }
else if (llval >= UCHAR_MAX) {
819 return static_cast<unsigned char>(fval);
823 __FP8_HOST_DEVICE__
operator unsigned int()
const {
824 if (internal::hip_fp8_fnuz_is_nan(__x)) {
829 auto llval =
static_cast<long long>(fval);
833 return static_cast<unsigned int>(fval);
837 __FP8_HOST_DEVICE__
operator unsigned long int()
const {
838 if (internal::hip_fp8_fnuz_is_nan(__x)) {
843 auto llval =
static_cast<long long>(fval);
847 return static_cast<unsigned long>(fval);
851 __FP8_HOST_DEVICE__
operator unsigned long long int()
const {
852 if (internal::hip_fp8_fnuz_is_nan(__x)) {
857 auto llval =
static_cast<long long>(fval);
861 return static_cast<unsigned long long>(fval);
865 __FP8_HOST_DEVICE__
operator unsigned short int()
const {
866 if (internal::hip_fp8_fnuz_is_nan(__x)) {
871 auto llval =
static_cast<long long>(fval);
875 return static_cast<unsigned short>(fval);
887 static constexpr unsigned int __we = 4;
888 static constexpr unsigned int __wm = 3;
910 __FP8_HOST_DEVICE__
operator __half2()
const {
915 __FP8_HOST_DEVICE__
operator float2()
const {
916#if HIP_FP8_CVT_FAST_PATH
917 return internal::cast_to_f32x2_from_f8x2(__x, __default_interpret);
935 static constexpr unsigned int __we = 4;
936 static constexpr unsigned int __wm = 3;
971 reinterpret_cast<unsigned short>(
973 reinterpret_cast<unsigned short>(
990 __FP8_HOST_DEVICE__
operator float4()
const {
994#if HIP_FP8_CVT_FAST_PATH
995 float2 high = internal::cast_to_f32x2_from_f8x2(fp8x2_high, __default_interpret);
996 float2 low = internal::cast_to_f32x2_from_f8x2(fp8x2_low, __default_interpret);
998 float2 high =
float2(internal::cast_from_f8<float, true>(
1000 internal::cast_from_f8<float, true>(
1002 float2 low =
float2(internal::cast_from_f8<float, true>(
1004 internal::cast_from_f8<float, true>(
1007 return float4(low.x, low.y, high.x, high.y);
1019 static constexpr unsigned int __we = 5;
1020 static constexpr unsigned int __wm = 2;
1029 __default_interpret)) {}
1034 __default_interpret)) {}
1039 __default_interpret)) {}
1044 __default_interpret)) {}
1049 __default_interpret)) {}
1054 __default_interpret)) {}
1067 __default_interpret)) {}
1072 __default_interpret)) {}
1078 __FP8_HOST_DEVICE__
operator float()
const {
1079#if HIP_FP8_CVT_FAST_PATH
1080 return internal::cast_to_f32_from_f8(__x, __default_interpret);
1082 return internal::cast_from_f8<float, true>(__x, __wm, __we);
1087 __FP8_HOST_DEVICE__
operator __half()
const {
1092 __FP8_HOST_DEVICE__
operator __hip_bfloat16()
const {
1094 return __hip_bfloat16(f);
1098 __FP8_HOST_DEVICE__
operator bool()
const {
1100 return !(
static_cast<unsigned short>(__x) == 0);
1104 __FP8_HOST_DEVICE__
operator char()
const {
1105 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1110 auto llval =
static_cast<long long>(fval);
1111 if (llval <= CHAR_MIN) {
1113 }
else if (llval >= CHAR_MAX) {
1116 return static_cast<char>(fval);
1120 __FP8_HOST_DEVICE__
operator double()
const {
1121 return internal::cast_from_f8<double, true>(__x, __wm, __we);
1125 __FP8_HOST_DEVICE__
operator int()
const {
1126 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1131 return static_cast<int>(fval);
1135 __FP8_HOST_DEVICE__
operator long int()
const {
1136 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1141 return static_cast<long>(fval);
1145 __FP8_HOST_DEVICE__
operator long long int()
const {
1146 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1151 return static_cast<long long>(fval);
1155 __FP8_HOST_DEVICE__
operator short int()
const {
1156 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1161 auto llval =
static_cast<long long>(fval);
1162 if (llval <= SHRT_MIN) {
1164 }
else if (llval >= SHRT_MAX) {
1167 return static_cast<short>(fval);
1171 __FP8_HOST_DEVICE__
operator signed char()
const {
1172 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1177 auto llval =
static_cast<long long>(fval);
1178 if (llval <= SCHAR_MIN) {
1180 }
else if (llval >= SCHAR_MAX) {
1183 return static_cast<signed char>(fval);
1187 __FP8_HOST_DEVICE__
operator unsigned char()
const {
1188 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1193 auto llval =
static_cast<long long>(fval);
1196 }
else if (llval >= UCHAR_MAX) {
1199 return static_cast<unsigned char>(fval);
1203 __FP8_HOST_DEVICE__
operator unsigned int()
const {
1204 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1209 auto llval =
static_cast<long long>(fval);
1213 return static_cast<unsigned int>(fval);
1217 __FP8_HOST_DEVICE__
operator unsigned long int()
const {
1218 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1223 auto llval =
static_cast<long long>(fval);
1227 return static_cast<unsigned long>(fval);
1231 __FP8_HOST_DEVICE__
operator unsigned long long int()
const {
1232 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1237 auto llval =
static_cast<long long>(fval);
1241 return static_cast<unsigned long long>(fval);
1245 __FP8_HOST_DEVICE__
operator unsigned short int()
const {
1246 if (internal::hip_fp8_fnuz_is_nan(__x)) {
1251 auto llval =
static_cast<long long>(fval);
1255 return static_cast<unsigned short>(fval);
1267 static constexpr unsigned int __we = 5;
1268 static constexpr unsigned int __wm = 2;
1290 __FP8_HOST_DEVICE__
operator __half2()
const {
1295 __FP8_HOST_DEVICE__
operator float2()
const {
1296#if HIP_FP8_CVT_FAST_PATH
1297 return internal::cast_to_f32x2_from_f8x2(__x, __default_interpret);
1315 static constexpr unsigned int __we = 5;
1316 static constexpr unsigned int __wm = 2;
1351 reinterpret_cast<unsigned short>(
1353 reinterpret_cast<unsigned short>(
1370 __FP8_HOST_DEVICE__
operator float4()
const {
1374#if HIP_FP8_CVT_FAST_PATH
1375 float2 high = internal::cast_to_f32x2_from_f8x2(fp8x2_high, __default_interpret);
1376 float2 low = internal::cast_to_f32x2_from_f8x2(fp8x2_low, __default_interpret);
1378 float2 high =
float2(internal::cast_from_f8<float, true>(
1380 internal::cast_from_f8<float, true>(
1382 float2 low =
float2(internal::cast_from_f8<float, true>(
1384 internal::cast_from_f8<float, true>(
1387 return float4(low.x, low.y, high.x, high.y);
__hip_saturation_t
Describes saturation behavior.
Definition amd_hip_fp8.h:76
@ __HIP_SATFINITE
Definition amd_hip_fp8.h:78
@ __HIP_NOSAT
Definition amd_hip_fp8.h:77
__FP8_HOST_DEVICE_STATIC__ __half2_raw __hip_cvt_fp8x2_to_halfraw2(const __hip_fp8x2_storage_t x, const __hip_fp8_interpretation_t type)
convert __hip_fp8x2_storage_t to __half2_raw
Definition amd_hip_fp8.h:598
__hip_fp8_interpretation_t
Describes FP8 interpretation.
Definition amd_hip_fp8.h:68
@ __HIP_E4M3_FNUZ
Definition amd_hip_fp8.h:69
@ __HIP_E5M2_FNUZ
Definition amd_hip_fp8.h:70
__FP8_HOST_DEVICE_STATIC__ __hip_fp8x2_storage_t __hip_cvt_double2_to_fp8x2(const double2 d2, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert double2 to __hip_fp8x2_storage_t
Definition amd_hip_fp8.h:539
__FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t __hip_cvt_double_to_fp8(const double d, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert double to __hip_fp8_storage_t
Definition amd_hip_fp8.h:524
__FP8_HOST_DEVICE_STATIC__ __hip_fp8x2_storage_t __hip_cvt_halfraw2_to_fp8x2(const __half2_raw x, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert __half2_raw to __hip_fp8x2_storage_t
Definition amd_hip_fp8.h:627
unsigned short int __hip_fp8x2_storage_t
type to store two fp8 numbers
Definition amd_hip_fp8.h:92
__FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t __hip_cvt_halfraw_to_fp8(const __half_raw x, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert __half_raw to __hip_fp8_storage_t
Definition amd_hip_fp8.h:614
__FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t __hip_cvt_float_to_fp8(const float f, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert float to __hip_fp8_storage_t
Definition amd_hip_fp8.h:486
__FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t __hip_cvt_bfloat16raw_to_fp8(const __hip_bfloat16_raw hr, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert __hip_bfloat16_raw to __hip_fp8_storage_t
Definition amd_hip_fp8.h:555
unsigned int __hip_fp8x4_storage_t
type to store four fp8 numbers
Definition amd_hip_fp8.h:99
__FP8_HOST_DEVICE_STATIC__ __hip_fp8x2_storage_t __hip_cvt_bfloat16raw2_to_fp8x2(const __hip_bfloat162_raw hr, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert double2 to __hip_fp8x2_storage_t
Definition amd_hip_fp8.h:570
unsigned char __hip_fp8_storage_t
type to store single fp8 number
Definition amd_hip_fp8.h:85
__FP8_HOST_DEVICE_STATIC__ __half_raw __hip_cvt_fp8_to_halfraw(const __hip_fp8_storage_t x, const __hip_fp8_interpretation_t type)
convert __hip_fp8_storage_t to __half_raw
Definition amd_hip_fp8.h:584
__FP8_HOST_DEVICE_STATIC__ __hip_fp8x2_storage_t __hip_cvt_float2_to_fp8x2(const float2 f2, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type)
convert float2 to __hip_fp8x2_storage_t
Definition amd_hip_fp8.h:505
hip_bf16.h provides struct for __hip_bfloat16 types
struct representing single fp8 number with e4m3 interpretation
Definition amd_hip_fp8.h:636
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const unsigned int val)
Definition amd_hip_fp8.h:667
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz()=default
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const double f)
Definition amd_hip_fp8.h:677
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const __half f)
Definition amd_hip_fp8.h:690
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const short int val)
Definition amd_hip_fp8.h:657
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const float f)
Definition amd_hip_fp8.h:681
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const __hip_bfloat16 f)
Definition amd_hip_fp8.h:685
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const long int val)
Definition amd_hip_fp8.h:647
static constexpr __hip_saturation_t __default_saturation
raw storage of fp8 number
Definition amd_hip_fp8.h:638
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const unsigned long int val)
Definition amd_hip_fp8.h:662
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const unsigned short int val)
Definition amd_hip_fp8.h:672
__FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const int val)
Definition amd_hip_fp8.h:652
struct representing two fp8 numbers with e4m3 interpretation
Definition amd_hip_fp8.h:883
__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const float2 val)
Definition amd_hip_fp8.h:895
__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const double2 val)
Definition amd_hip_fp8.h:891
static constexpr __hip_saturation_t __default_saturation
raw storage of two fp8 numbers
Definition amd_hip_fp8.h:885
__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const __half2 val)
Definition amd_hip_fp8.h:903
__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz()=default
__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const __hip_bfloat162 val)
Definition amd_hip_fp8.h:899
struct representing four fp8 numbers with e4m3 interpretation
Definition amd_hip_fp8.h:931
__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const __hip_bfloat162 low, const __hip_bfloat162 high)
Definition amd_hip_fp8.h:969
__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz()=default
__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const double4 val)
Definition amd_hip_fp8.h:939
static constexpr __hip_saturation_t __default_saturation
raw storage of four fp8 numbers
Definition amd_hip_fp8.h:933
__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const __half2 low, const __half2 high)
Definition amd_hip_fp8.h:978
__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const float4 val)
Definition amd_hip_fp8.h:954
struct representing one fp8 number with e5m2 interpretation
Definition amd_hip_fp8.h:1015
static constexpr __hip_saturation_t __default_saturation
raw storage of one fp8 numbers
Definition amd_hip_fp8.h:1017
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const unsigned short int val)
Definition amd_hip_fp8.h:1052
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz()=default
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const __hip_bfloat16 f)
Definition amd_hip_fp8.h:1065
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const long int val)
Definition amd_hip_fp8.h:1027
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const unsigned int val)
Definition amd_hip_fp8.h:1047
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const __half f)
Definition amd_hip_fp8.h:1070
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const int val)
Definition amd_hip_fp8.h:1032
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const double f)
Definition amd_hip_fp8.h:1057
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const short int val)
Definition amd_hip_fp8.h:1037
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const float f)
Definition amd_hip_fp8.h:1061
__FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const unsigned long int val)
Definition amd_hip_fp8.h:1042
struct representing two fp8 numbers with e5m2 interpretation
Definition amd_hip_fp8.h:1263
__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const float2 val)
Definition amd_hip_fp8.h:1275
static constexpr __hip_saturation_t __default_saturation
raw storage of two fp8 numbers
Definition amd_hip_fp8.h:1265
__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const __half2 val)
Definition amd_hip_fp8.h:1283
__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const __hip_bfloat162 val)
Definition amd_hip_fp8.h:1279
__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const double2 val)
Definition amd_hip_fp8.h:1271
__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz()=default
struct representing four fp8 numbers with e5m2 interpretation
Definition amd_hip_fp8.h:1311
__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const __hip_bfloat162 low, const __hip_bfloat162 high)
Definition amd_hip_fp8.h:1349
__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const float4 val)
Definition amd_hip_fp8.h:1334
static constexpr __hip_saturation_t __default_saturation
raw storage of four fp8 numbers
Definition amd_hip_fp8.h:1313
__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const __half2 low, const __half2 high)
Definition amd_hip_fp8.h:1358
__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const double4 val)
Definition amd_hip_fp8.h:1319
Definition amd_hip_vector_types.h:2035
Definition amd_hip_vector_types.h:2042
Definition amd_hip_vector_types.h:2072
Definition amd_hip_vector_types.h:2079
Definition hip_fp16_gcc.h:7
Definition hip_fp16_gcc.h:11