HIP: Heterogenous-computing Interface for Portability
|
amd_hip_fp8.h header, for AMD fp8 data types More...
#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 "amd_hip_bf16.h"
#include "math_fwd.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... | |
Macros | |
#define | HIP_FP8_CVT_FAST_PATH 0 |
#define | __FP8_HOST_DEVICE__ __host__ __device__ |
#define | __FP8_HOST_DEVICE_STATIC__ __FP8_HOST_DEVICE__ static inline |
Enumerations | |
enum | __hip_fp8_interpretation_t { __HIP_E4M3_FNUZ = 0 , __HIP_E5M2_FNUZ = 1 } |
Describes FP8 interpretation. More... | |
enum | __hip_saturation_t { __HIP_NOSAT = 0 , __HIP_SATFINITE = 1 } |
Describes saturation behavior. More... | |
Functions | |
template<typename T , bool negative_zero_nan> | |
__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 negative_zero_nan> | |
__FP8_HOST_DEVICE_STATIC__ T | internal::cast_from_f8 (__hip_fp8_storage_t x, int wm, int we) |
__FP8_HOST_DEVICE_STATIC__ bool | internal::hip_fp8_fnuz_is_nan (__hip_fp8_storage_t a) |
__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 | |
__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 | |
__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 | |
__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 | |
__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 | |
__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 | |
__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 | |
__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 | |
__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 | |
__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 | |
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.
enum __hip_saturation_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 | type ) |
convert double2 to __hip_fp8x2_storage_t
hr | __hip_bfloat162_raw value |
sat | saturation of fp8 |
type | interpretation of 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 | type ) |
convert __hip_bfloat16_raw to __hip_fp8_storage_t
hr | __hip_bfloat16_raw val |
sat | saturation of fp8 |
type | interpretation of fp8 |
__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 ) |
__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
d | double val |
sat | saturation of fp8 |
type | interpretation of fp8 |
__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 ) |
__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
f | float number |
sat | saturation of fp8 |
type | interpretation of fp8 |
__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
x | __hip_fp8_storage_t val |
type | interpretation of fp8 |
__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
x | __hip_fp8x2_storage_t val |
type | interpretation of fp8 |
__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
x | __half2_raw value |
sat | saturation of fp8 |
type | interpretation of 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 | type ) |
convert __half_raw to __hip_fp8_storage_t
x | __half_raw value |
sat | saturation of fp8 |
type | interpretation of fp8 |