e8m0_utils.hpp 894 Bytes
Newer Older
1
2
3
4
5
6
7
8
9
10
// 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 {

Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
11
__host__ __device__ inline float cast_to_float(e8m0_bexp_t const bexp)
12
{
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
13
    // TODO: check performance and try bit shift impl
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
14
    return std::powf(2, bit_cast<uint8_t>(bexp) - NumericUtils<e8m0_bexp_t>::bias);
15
16
}

Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
17
__host__ __device__ inline e8m0_bexp_t cast_from_float(float const scale)
18
{
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
19
20
    uint32_t e = bit_cast<uint32_t>(scale) & NumericUtils<float>::nan_mask;
    return static_cast<uint8_t>(e >> 23);
21
22
}

Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
23
template <>
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
24
__host__ __device__ inline int get_exponent_value<e8m0_bexp_t>(e8m0_bexp_t x)
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
25
{
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
26
    x.data >>= NumericUtils<e8m0_bexp_t>::mant;
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
27

Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
28
    x.data &= ((1 << NumericUtils<e8m0_bexp_t>::exp) - 1);
Rostyslav Geyyer's avatar
Rostyslav Geyyer committed
29
30
31
32

    return static_cast<int>(x.data);
}

33
} // namespace ck::utils