HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_hip_fp8.h File Reference

amd_hip_fp8.h header, for AMD fp8 data types More...

#include "amd_hip_bf16.h"
#include <hip/amd_detail/amd_hip_common.h>
#include <climits>
#include "host_defines.h"
#include "amd_hip_vector_types.h"
#include "amd_hip_fp16.h"
#include "math_fwd.h"
#include "hip_assert.h"

Go to the source code of this file.

Classes

struct  __hip_fp8_e4m3_fnuz
 struct representing single fp8 number with e4m3 interpretation More...
 
struct  __hip_fp8x2_e4m3_fnuz
 struct representing two fp8 numbers with e4m3 interpretation More...
 
struct  __hip_fp8x4_e4m3_fnuz
 struct representing four fp8 numbers with e4m3 interpretation More...
 
struct  __hip_fp8_e5m2_fnuz
 struct representing one fp8 number with e5m2 interpretation More...
 
struct  __hip_fp8x2_e5m2_fnuz
 struct representing two fp8 numbers with e5m2 interpretation More...
 
struct  __hip_fp8x4_e5m2_fnuz
 struct representing four fp8 numbers with e5m2 interpretation More...
 
struct  __hip_fp8_e4m3
 struct representing ocp fp8 numbers with e4m3 interpretation More...
 
struct  __hip_fp8x2_e4m3
 struct representing two ocp fp8 numbers with e4m3 interpretation More...
 
struct  __hip_fp8x4_e4m3
 struct representing four ocp fp8 numbers with e4m3 interpretation More...
 
struct  __hip_fp8_e5m2
 struct representing ocp fp8 numbers with e5m2 interpretation More...
 
struct  __hip_fp8x2_e5m2
 struct representing two ocp fp8 numbers with e5m2 interpretation More...
 
struct  __hip_fp8x4_e5m2
 struct representing four ocp fp8 numbers with e5m2 interpretation More...
 

Macros

#define HIP_FP8_CVT_FAST_PATH   0
 
#define HIP_FP8_TYPE_FNUZ   1
 
#define HIP_FP8_TYPE_OCP   1
 
#define __HIP_SCHAR_MAX   SCHAR_MAX
 
#define __HIP_SCHAR_MIN   SCHAR_MIN
 
#define __HIP_UCHAR_MAX   UCHAR_MAX
 
#define __HIP_SHRT_MIN   SHRT_MIN
 
#define __HIP_SHRT_MAX   SHRT_MAX
 
#define __HIP_CHAR_MIN   CHAR_MIN
 
#define __HIP_CHAR_MAX   CHAR_MAX
 
#define __FP8_HOST_DEVICE__   __host__ __device__
 
#define __FP8_HOST_DEVICE_STATIC__   __FP8_HOST_DEVICE__ static inline
 
#define __FP8_HOST__   __host__
 
#define __FP8_HOST_STATIC__   __FP8_HOST__ static inline
 
#define __assert_ocp_support(interp)
 
#define __assert_fnuz_support(interp)
 

Typedefs

typedef unsigned char __hip_fp8_storage_t
 type to store single fp8 number
 
typedef unsigned short int __hip_fp8x2_storage_t
 type to store two fp8 numbers
 
typedef unsigned int __hip_fp8x4_storage_t
 type to store four fp8 numbers
 

Enumerations

enum  __hip_fp8_interpretation_t { __HIP_E4M3 = 0 , __HIP_E5M2 = 1 , __HIP_E4M3_FNUZ = 2 , __HIP_E5M2_FNUZ = 3 }
 Describes FP8 interpretation. More...
 
enum  __hip_saturation_t { __HIP_NOSAT = 0 , __HIP_SATFINITE = 1 }
 Describes saturation behavior. More...
 

Functions

__FP8_HOST_DEVICE_STATIC__ void internal::__is_interpret_supported (__hip_fp8_interpretation_t interp)
 
template<typename T, bool is_fnuz>
__FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t internal::cast_to_f8 (T _x, int wm, int we, bool clip=false, bool stoch=false, unsigned int rng=0)
 
template<typename T, bool is_fnuz>
__FP8_HOST_DEVICE_STATIC__ T internal::cast_from_f8 (__hip_fp8_storage_t x, int wm, int we, bool clip=false)
 
__FP8_HOST_DEVICE_STATIC__ bool internal::hip_fp8_fnuz_is_nan (__hip_fp8_storage_t a)
 
__FP8_HOST_DEVICE_STATIC__ bool internal::hip_fp8_ocp_is_nan (__hip_fp8_storage_t a, const __hip_fp8_interpretation_t type)
 
__FP8_HOST_DEVICE_STATIC__ bool internal::hip_fp8_ocp_is_inf (__hip_fp8_storage_t a, const __hip_fp8_interpretation_t type)
 
__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 interp)
 convert float to __hip_fp8_storage_t
 
__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 interp)
 convert float2 to __hip_fp8x2_storage_t
 
__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 interp)
 convert double to __hip_fp8_storage_t
 
__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 interp)
 convert double2 to __hip_fp8x2_storage_t
 
__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 interp)
 convert __hip_bfloat16_raw to __hip_fp8_storage_t
 
__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 interp)
 convert double2 to __hip_fp8x2_storage_t
 
__FP8_HOST_DEVICE_STATIC__ __half_raw __hip_cvt_fp8_to_halfraw (const __hip_fp8_storage_t x, const __hip_fp8_interpretation_t interp)
 convert __hip_fp8_storage_t to __half_raw
 
__FP8_HOST_DEVICE_STATIC__ __half2_raw __hip_cvt_fp8x2_to_halfraw2 (const __hip_fp8x2_storage_t x, const __hip_fp8_interpretation_t interp)
 convert __hip_fp8x2_storage_t to __half2_raw
 
__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 interp)
 convert __half_raw to __hip_fp8_storage_t
 
__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 interp)
 convert __half2_raw to __hip_fp8x2_storage_t
 

Detailed Description

amd_hip_fp8.h header, for AMD fp8 data types

MIT License

Copyright (c) 2019 - 2024 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.

Macro Definition Documentation

◆ __assert_fnuz_support

#define __assert_fnuz_support ( interp)
Value:
{ \
if (interp != __HIP_E4M3_FNUZ && interp != __HIP_E5M2_FNUZ) { \
__hip_assert(false && "type is unsupported by current target device"); \
} \
}
@ __HIP_E4M3_FNUZ
Definition amd_hip_fp8.h:125
@ __HIP_E5M2_FNUZ
Definition amd_hip_fp8.h:126

◆ __assert_ocp_support

#define __assert_ocp_support ( interp)
Value:
{ \
if (interp != __HIP_E4M3 && interp != __HIP_E5M2) { \
__hip_assert(false && "type is unsupported by current target device"); \
} \
}
@ __HIP_E5M2
Definition amd_hip_fp8.h:124
@ __HIP_E4M3
Definition amd_hip_fp8.h:123

Enumeration Type Documentation

◆ __hip_fp8_interpretation_t

Describes FP8 interpretation.

Enumerator
__HIP_E4M3 

OCP E4M3

__HIP_E5M2 

OCP E5M2

__HIP_E4M3_FNUZ 

Standard FP8

__HIP_E5M2_FNUZ 

BF8

◆ __hip_saturation_t

Describes saturation behavior.

Enumerator
__HIP_NOSAT 

No saturation

__HIP_SATFINITE 

Saturate to finite

Function Documentation

◆ __hip_cvt_bfloat16raw2_to_fp8x2()

__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 interp )

convert double2 to __hip_fp8x2_storage_t

Parameters
hr__hip_bfloat162_raw value
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8x2_storage_t

◆ __hip_cvt_bfloat16raw_to_fp8()

__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 interp )

convert __hip_bfloat16_raw to __hip_fp8_storage_t

Parameters
hr__hip_bfloat16_raw val
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8_storage_t

◆ __hip_cvt_double2_to_fp8x2()

__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 interp )

convert double2 to __hip_fp8x2_storage_t

Parameters
d2double2 val
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8x2_storage_t

◆ __hip_cvt_double_to_fp8()

__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 interp )

convert double to __hip_fp8_storage_t

Parameters
ddouble val
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8_storage_t

◆ __hip_cvt_float2_to_fp8x2()

__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 interp )

convert float2 to __hip_fp8x2_storage_t

Parameters
f2float2 number
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8x2_storage_t

◆ __hip_cvt_float_to_fp8()

__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 interp )

convert float to __hip_fp8_storage_t

Parameters
ffloat number
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8_storage_t

◆ __hip_cvt_fp8_to_halfraw()

__FP8_HOST_DEVICE_STATIC__ __half_raw __hip_cvt_fp8_to_halfraw ( const __hip_fp8_storage_t x,
const __hip_fp8_interpretation_t interp )

convert __hip_fp8_storage_t to __half_raw

Parameters
x__hip_fp8_storage_t val
interpinterpretation of fp8
Returns
__half_raw

◆ __hip_cvt_fp8x2_to_halfraw2()

__FP8_HOST_DEVICE_STATIC__ __half2_raw __hip_cvt_fp8x2_to_halfraw2 ( const __hip_fp8x2_storage_t x,
const __hip_fp8_interpretation_t interp )

convert __hip_fp8x2_storage_t to __half2_raw

Parameters
x__hip_fp8x2_storage_t val
interpinterpretation of fp8
Returns
__half2_raw

◆ __hip_cvt_halfraw2_to_fp8x2()

__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 interp )

convert __half2_raw to __hip_fp8x2_storage_t

Parameters
x__half2_raw value
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8x2_storage_t

◆ __hip_cvt_halfraw_to_fp8()

__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 interp )

convert __half_raw to __hip_fp8_storage_t

Parameters
x__half_raw value
satsaturation of fp8
interpinterpretation of fp8
Returns
__hip_fp8_storage_t