/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/amd_ck_fp8.hpp File Reference

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/amd_ck_fp8.hpp File Reference#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/amd_ck_fp8.hpp File Reference
Classes | Namespaces | Macros | Typedefs | Enumerations | Functions
amd_ck_fp8.hpp File Reference
#include "ck/ck.hpp"
#include "ck/utility/enable_if.hpp"
#include "ck/utility/get_id.hpp"
#include "ck/utility/random_gen.hpp"
#include "ck/utility/functional.hpp"
#include "ck/utility/type.hpp"

Go to the source code of this file.

Classes

struct  ck::f8_ocp_t
 
struct  ck::bf8_ocp_t
 

Namespaces

 ck
 
 ck::fp8_impl
 

Macros

#define CK_USE_FNUZ_FP8   0
 
#define CK_USE_OCP_FP8   0
 
#define CK_FP8_CVT_FAST_PATH   0
 
#define CK_OCP_FP8_CVT_FAST_PATH   0
 
#define __fp8_impl_assert_ocp_support(interp)
 
#define __fp8_impl_assert_fnuz_support(interp)
 
#define CK_FP8_TYPE_FNUZ   1
 
#define CK_FP8_TYPE_OCP   0
 

Typedefs

using ck::f8_fnuz_t = _BitInt(8)
 
using ck::bf8_fnuz_t = unsigned _BitInt(8)
 
typedef unsigned char ck::fp8_storage_t
 
typedef fp8_storage_t ck::fp8_impl::fp8x2_storage_t
 
typedef _Float16 ck::fp8_impl::half2_t
 
typedef ushort ck::fp8_impl::ushortx2_t
 
typedef short ck::fp8_impl::shortx2_t
 
typedef float ck::fp8_impl::float2_t
 
using ck::f8_t = f8_fnuz_t
 
using ck::bf8_t = bf8_fnuz_t
 

Enumerations

enum class  ck::ck_fp8_interpretation_t {
  ck::CK_E4M3_OCP = 0 ,
  ck::CK_E5M2_OCP = 1 ,
  ck::CK_E4M3_FNUZ = 2 ,
  ck::CK_E5M2_FNUZ = 3
}
 Describes FP8 interpretation. More...
 
enum class  ck::ck_saturation_t {
  ck::CK_NOSAT = 0 ,
  ck::CK_SATFINITE = 1
}
 Describes saturation behavior. More...
 

Functions

template<>
__host__ constexpr __device__ bool ck::fp8_is_nan (f8_ocp_t a)
 
template<>
__host__ constexpr __device__ bool ck::fp8_is_nan (bf8_ocp_t a)
 
template<>
__host__ constexpr __device__ bool ck::fp8_is_nan (f8_fnuz_t a)
 
template<>
__host__ constexpr __device__ bool ck::fp8_is_nan (bf8_fnuz_t a)
 
template<>
__host__ constexpr __device__ bool ck::fp8_is_inf (bf8_ocp_t a)
 

Macro Definition Documentation

◆ __fp8_impl_assert_fnuz_support

#define __fp8_impl_assert_fnuz_support (   interp)
Value:
{ \
if(interp != ck_fp8_interpretation_t::CK_E4M3_FNUZ && \
interp != ck_fp8_interpretation_t::CK_E5M2_FNUZ) \
{ \
__hip_assert(false && "type is unsupported by current target device"); \
} \
}

◆ __fp8_impl_assert_ocp_support

#define __fp8_impl_assert_ocp_support (   interp)
Value:
{ \
if(interp != ck_fp8_interpretation_t::CK_E4M3_OCP && \
interp != ck_fp8_interpretation_t::CK_E5M2_OCP) \
{ \
__hip_assert(false && "type is unsupported by current target device"); \
} \
}

◆ CK_FP8_CVT_FAST_PATH

#define CK_FP8_CVT_FAST_PATH   0

◆ CK_FP8_TYPE_FNUZ

#define CK_FP8_TYPE_FNUZ   1

◆ CK_FP8_TYPE_OCP

#define CK_FP8_TYPE_OCP   0

◆ CK_OCP_FP8_CVT_FAST_PATH

#define CK_OCP_FP8_CVT_FAST_PATH   0

◆ CK_USE_FNUZ_FP8

#define CK_USE_FNUZ_FP8   0

◆ CK_USE_OCP_FP8

#define CK_USE_OCP_FP8   0