e8m0.hpp Source File

e8m0.hpp Source File#

Composable Kernel: e8m0.hpp Source File
// SPDX-License-Identifier: MIT
// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
 
#pragma once
 
#ifndef CK_CODE_GEN_RTC
#include "ck/utility/type.hpp"
 
namespace ck {
struct e8m0_bexp_t
{
    using type = uint8_t;
    type data;
 
    constexpr static type bias     = 127;
    constexpr static type nan_mask = 0xFF;
 
    __host__ __device__ constexpr e8m0_bexp_t() : data{type{}} {}
    __host__ __device__ constexpr e8m0_bexp_t(type init) : data{init} {}
    __host__ __device__ constexpr e8m0_bexp_t(int init) : data{static_cast<type>(init & nan_mask)}
    {
    }
    __host__ __device__ explicit constexpr e8m0_bexp_t(float scale)
        : data{static_cast<type>((bit_cast<uint32_t>(scale) & (nan_mask << 23)) >> 23)}
    {
    }
 
    __host__ __device__ explicit constexpr operator float() const
    {
        if(data == nan_mask || data == 0)
        {
            uint32_t bits = data << 1;
            bits |= 1;
            bits <<= 22;
            return bit_cast<float>(bits);
        }
        else
        {
            uint32_t bits = data << 23;
            return bit_cast<float>(bits);
        }
    }
 
    __host__ __device__ constexpr bool operator==(const e8m0_bexp_t& other) const
    {
        // strict IEEE compliance for NaN
        return data == other.data && data != nan_mask;
    }
 
    __host__ __device__ constexpr bool is_nan() const { return data == nan_mask; }
};
 
namespace utils {
 
template <typename T>
__host__ __device__ inline constexpr int32_t get_exponent_value(T x);
 
template <>
__host__ __device__ inline constexpr int32_t get_exponent_value<e8m0_bexp_t>(e8m0_bexp_t x)
{
    return x.data;
}
 
} // namespace utils
 
} // namespace ck
#endif