// SPDX-License-Identifier: MIT // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once #include "ck/utility/data_type.hpp" #include "ck/utility/mxfp_utils.hpp" namespace ck::utils { __host__ __device__ inline float cast_to_float(e8m0_scale_t const scale) { // TODO: check performance and try bit shift impl return std::powf(2, bit_cast(scale) - NumericUtils::bias); } __host__ __device__ inline e8m0_scale_t cast_from_float(float const scale) { uint32_t e = bit_cast(scale) & NumericUtils::nan_mask; return static_cast(e >> 23); } template <> __host__ __device__ inline int get_exponent_value(e8m0_scale_t x) { x.data >>= NumericUtils::mant; x.data &= ((1 << NumericUtils::exp) - 1); return static_cast(x.data); } } // namespace ck::utils