amd_ck_fp8.hpp File Reference

amd_ck_fp8.hpp File Reference#

Composable Kernel: 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_fnuz_t
struct  ck::bf8_fnuz_t
struct  ck::f8_ocp_t
struct  ck::bf8_ocp_t

Namespaces

namespace  ck
namespace  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

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__ __device__ constexpr bool ck::fp8_is_nan (f8_ocp_t a)
template<>
__host__ __device__ constexpr bool ck::fp8_is_nan (bf8_ocp_t a)
template<>
__host__ __device__ constexpr bool ck::fp8_is_nan (f8_fnuz_t a)
template<>
__host__ __device__ constexpr bool ck::fp8_is_nan (bf8_fnuz_t a)
template<>
__host__ __device__ constexpr 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