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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/e8m0.hpp Source File
e8m0.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #ifndef CK_CODE_GEN_RTC
7 #include "ck/utility/type.hpp"
8 
9 namespace ck {
10 
25 struct e8m0_bexp_t
26 {
27  using type = uint8_t;
28  type data;
29 
30  constexpr static type bias = 127;
31  constexpr static type nan_mask = 0xFF;
32 
33  __host__ __device__ constexpr e8m0_bexp_t() : data{type{}} {}
34  __host__ __device__ constexpr e8m0_bexp_t(type init) : data{init} {}
35  __host__ __device__ constexpr e8m0_bexp_t(int init) : data{static_cast<type>(init & nan_mask)}
36  {
37  }
38  __host__ __device__ explicit constexpr e8m0_bexp_t(float scale)
39  : data{static_cast<type>((bit_cast<uint32_t>(scale) & (nan_mask << 23)) >> 23)}
40  {
41  }
42 
43  __host__ __device__ explicit constexpr operator float() const
44  {
45  if(data == nan_mask || data == 0)
46  {
47  uint32_t bits = data << 1;
48  bits |= 1;
49  bits <<= 22;
50  return bit_cast<float>(bits);
51  }
52  else
53  {
54  uint32_t bits = data << 23;
55  return bit_cast<float>(bits);
56  }
57  }
58 
59  __host__ __device__ constexpr bool operator==(const e8m0_bexp_t& other) const
60  {
61  // strict IEEE compliance for NaN
62  return data == other.data && data != nan_mask;
63  }
64 
65  __host__ __device__ constexpr bool is_nan() const { return data == nan_mask; }
66 };
67 
68 namespace utils {
69 
70 template <typename T>
71 __host__ __device__ inline constexpr int32_t get_exponent_value(T x);
72 
73 template <>
74 __host__ __device__ inline constexpr int32_t get_exponent_value<e8m0_bexp_t>(e8m0_bexp_t x)
75 {
76  return x.data;
77 }
78 
79 } // namespace utils
80 
81 } // namespace ck
82 #endif
ck::utils::get_exponent_value< e8m0_bexp_t >
__host__ constexpr __device__ int32_t get_exponent_value< e8m0_bexp_t >(e8m0_bexp_t x)
Definition: e8m0.hpp:74
ck::utils::get_exponent_value
__host__ constexpr __device__ int32_t get_exponent_value(T x)
Definition: mxfp_utils.hpp:35
ck
Definition: ck.hpp:268
ck::bit_cast
__host__ constexpr __device__ Y bit_cast(const X &x)
Definition: type.hpp:306
uint32_t
unsigned int uint32_t
Definition: stdint.h:126
int32_t
signed int int32_t
Definition: stdint.h:123
uint8_t
unsigned char uint8_t
Definition: stdint.h:124
ck::e8m0_bexp_t
Unsigned representation of a conventional biased Float32 exponent.
Definition: e8m0.hpp:26
ck::e8m0_bexp_t::is_nan
__host__ constexpr __device__ bool is_nan() const
Definition: e8m0.hpp:65
ck::e8m0_bexp_t::bias
constexpr static type bias
Definition: e8m0.hpp:30
ck::e8m0_bexp_t::data
type data
Definition: e8m0.hpp:28
ck::e8m0_bexp_t::type
uint8_t type
Definition: e8m0.hpp:27
ck::e8m0_bexp_t::e8m0_bexp_t
__host__ constexpr __device__ e8m0_bexp_t(type init)
Definition: e8m0.hpp:34
ck::e8m0_bexp_t::e8m0_bexp_t
__host__ constexpr __device__ e8m0_bexp_t(float scale)
Definition: e8m0.hpp:38
ck::e8m0_bexp_t::nan_mask
constexpr static type nan_mask
Definition: e8m0.hpp:31
ck::e8m0_bexp_t::e8m0_bexp_t
__host__ constexpr __device__ e8m0_bexp_t()
Definition: e8m0.hpp:33
ck::e8m0_bexp_t::e8m0_bexp_t
__host__ constexpr __device__ e8m0_bexp_t(int init)
Definition: e8m0.hpp:35
ck::e8m0_bexp_t::operator==
__host__ constexpr __device__ bool operator==(const e8m0_bexp_t &other) const
Definition: e8m0.hpp:59