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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/mxf6_utils.hpp Source File
Go to the documentation of this file.
 // SPDX-License-Identifier: MIT
 // Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
  
 #ifndef CK_CODE_GEN_RTC
 #pragma once
  
 #include "ck/utility/numeric_limits.hpp"
 #include "ck/utility/mxfp_utils.hpp"
  
 namespace ck::utils {
  
 template <>
 __host__ __device__ inline bool is_nan<f6_t>(e8m0_bexp_t const scale,
                                              f6_t const dataBytes [[maybe_unused]])
 {
     // no need to check for data as it does not have NaN representation
     return scale.is_nan();
 }
  
 template <>
 __host__ __device__ inline bool is_nan<bf6_t>(e8m0_bexp_t const scale,
                                               bf6_t const dataBytes [[maybe_unused]])
 {
     // no need to check for data as it does not have NaN representation
     return scale.is_nan();
 }
  
 template <>
 __host__ __device__ inline bool is_inf<f6_t>(e8m0_bexp_t const scale [[maybe_unused]],
                                              f6_t const data [[maybe_unused]])
 {
     // no inf representation for fp6
     return false;
 }
  
 template <>
 __host__ __device__ inline bool is_inf<bf6_t>(e8m0_bexp_t const scale [[maybe_unused]],
                                               bf6_t const data [[maybe_unused]])
 {
     // no inf representation for bf6
     return false;
 }
  
 template <>
 __host__ __device__ inline bool is_zero<f6_t>(e8m0_bexp_t const scale, f6_t const data)
 {
     if(is_nan<f6_t>(scale, data))
         return false;
  
     // no need to check for scale as it does not have a 0 representation
     f6_t result = (data & 0b00111111) & NumericUtils<f6_t>::set_sign_mask;
  
     return result == 0b0;
 }
  
 template <>
 __host__ __device__ inline bool is_zero<bf6_t>(e8m0_bexp_t const scale, bf6_t const data)
 {
     if(is_nan<bf6_t>(scale, data))
         return false;
  
     // no need to check for scale as it does not have a 0 representation
     bf6_t result = (data & 0b00111111) & NumericUtils<bf6_t>::set_sign_mask;
  
     return result == 0b0;
 }
  
 template <>
 __host__ __device__ inline float to_float<f6_t>(e8m0_bexp_t const scale, f6_t const data)
 {
     if(is_nan<f6_t>(scale, data))
         return NumericLimits<float>::QuietNaN();
  
     if(is_zero<f6_t>(scale, data))
         return 0.0f;
  
     f6_t prepared_data = data & 0b00111111;
  
     int scale_exp = get_exponent_value<e8m0_bexp_t>(scale);
  
     return convert_to_float<f6_t>(prepared_data, scale_exp);
 }
  
 template <>
 __host__ __device__ inline float to_float<bf6_t>(e8m0_bexp_t const scale, bf6_t const data)
 {
     if(is_nan<bf6_t>(scale, data))
         return NumericLimits<float>::QuietNaN();
  
     if(is_zero<bf6_t>(scale, data))
         return 0.0f;
  
     bf6_t prepared_data = data & 0b00111111;
  
     int scale_exp = get_exponent_value<e8m0_bexp_t>(scale);
  
     return convert_to_float<bf6_t>(prepared_data, scale_exp);
 }
  
 template <>
 __host__ __device__ inline f6_t sat_convert_to_type<f6_t>(float value)
 {
     cvt t;
     t.value_float = value;
     uint32_t sign = t.value_bitwise >> 31;
  
     if(std::isnan(value))
     {
  
         return sign ? NumericUtils<f6_t>::data_max_negative_normal_mask
                     : NumericUtils<f6_t>::data_max_positive_normal_mask;
     }
  
     if(std::abs(value) > NumericLimits<f6_t>::DataMaxNorm()) // covers inf case as well
         return sign ? NumericUtils<f6_t>::data_max_negative_normal_mask
                     : NumericUtils<f6_t>::data_max_positive_normal_mask;
  
     f6_t res = convert_to_type<f6_t>(value);
  
     if(std::abs(to_float<f6_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), res)) <
        NumericLimits<f6_t>::DataMinSubnorm())
         return sign ? NumericUtils<f6_t>::negative_zero_mask
                     : NumericUtils<f6_t>::positive_zero_mask;
  
     return res;
 }
  
 template <>
 __host__ __device__ inline bf6_t sat_convert_to_type<bf6_t>(float value)
 {
     cvt t;
     t.value_float = value;
     uint32_t sign = t.value_bitwise >> 31;
  
     if(std::isnan(value))
     {
  
         return sign ? NumericUtils<bf6_t>::data_max_negative_normal_mask
                     : NumericUtils<bf6_t>::data_max_positive_normal_mask;
     }
  
     if(std::abs(value) > NumericLimits<bf6_t>::DataMaxNorm()) // covers inf case as well
         return sign ? NumericUtils<bf6_t>::data_max_negative_normal_mask
                     : NumericUtils<bf6_t>::data_max_positive_normal_mask;
  
     bf6_t res = convert_to_type<bf6_t>(value);
  
     if(std::abs(to_float<bf6_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), res)) <
        NumericLimits<bf6_t>::DataMinSubnorm())
         return sign ? NumericUtils<bf6_t>::negative_zero_mask
                     : NumericUtils<bf6_t>::positive_zero_mask;
  
     return res;
 }
  
 template <>
 __host__ __device__ inline f6_t sat_convert_to_type_sr<f6_t>(float value, uint32_t seed)
 {
     cvt t;
     t.value_float = value;
     uint32_t sign = t.value_bitwise >> 31;
  
     if(std::isnan(value))
         return sign ? NumericUtils<f6_t>::data_max_negative_normal_mask
                     : NumericUtils<f6_t>::data_max_positive_normal_mask;
  
     if(std::abs(value) > NumericLimits<f6_t>::DataMaxNorm()) // covers inf case as well
         return sign ? NumericUtils<f6_t>::data_max_negative_normal_mask
                     : NumericUtils<f6_t>::data_max_positive_normal_mask;
  
     f6_t res = convert_to_type_sr<f6_t>(value, seed);
  
     if(std::abs(to_float<f6_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), res)) <
        NumericLimits<f6_t>::DataMinSubnorm())
         return sign ? NumericUtils<f6_t>::negative_zero_mask
                     : NumericUtils<f6_t>::positive_zero_mask;
  
     return res;
 }
  
 template <>
 __host__ __device__ inline bf6_t sat_convert_to_type_sr<bf6_t>(float value, uint32_t seed)
 {
     cvt t;
     t.value_float = value;
     uint32_t sign = t.value_bitwise >> 31;
  
     if(std::isnan(value))
         return sign ? NumericUtils<bf6_t>::data_max_negative_normal_mask
                     : NumericUtils<bf6_t>::data_max_positive_normal_mask;
     if(std::abs(value) > NumericLimits<bf6_t>::DataMaxNorm()) // covers inf case as well
         return sign ? NumericUtils<bf6_t>::data_max_negative_normal_mask
                     : NumericUtils<bf6_t>::data_max_positive_normal_mask;
  
     bf6_t res = convert_to_type_sr<bf6_t>(value, seed);
  
     if(std::abs(to_float<bf6_t>(NumericLimits<e8m0_bexp_t>::Binary_1(), res)) <
        NumericLimits<bf6_t>::DataMinSubnorm())
         return sign ? NumericUtils<bf6_t>::negative_zero_mask
                     : NumericUtils<bf6_t>::positive_zero_mask;
  
     return res;
 }
 } // namespace ck::utils
 #endif