type_convert.hpp File Reference#
type_convert.hpp File Reference
#include "ck/utility/data_type.hpp"#include "ck/utility/f8_utils.hpp"#include "ck/utility/get_id.hpp"#include "ck/utility/mxf4_utils.hpp"#include "ck/utility/mxf6_utils.hpp"#include "ck/utility/random_gen.hpp"#include "ck/utility/array.hpp"#include "ck/utility/amd_inline_asm.hpp"#include "ck/utility/type.hpp"Go to the source code of this file.
Namespaces | |
| namespace | ck |
| namespace | ck::details |
Functions | |
| template<typename Y, typename X> | |
| __host__ __device__ constexpr Y | ck::bf16_convert_rtn (X x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::bf16_convert_rtn< bhalf_t, float > (float x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::bf16_convert_rtn< bhalf_t, half_t > (half_t x) |
| template<typename Y, typename X, ck::enable_if_t<!(ck::is_const_v< Y >||ck::is_const_v< X >), bool > = false> | |
| __host__ __device__ constexpr Y | ck::type_convert (X x) |
| template<typename Y, typename X, ck::enable_if_t< ck::is_const_v< Y >||ck::is_const_v< X >, bool > = false> | |
| __host__ __device__ constexpr Y | ck::type_convert (X x) |
| template<> | |
| __host__ __device__ constexpr float | ck::type_convert< float, bhalf_t > (bhalf_t x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::type_convert< bhalf_t, float > (float x) |
| template<> | |
| __host__ __device__ constexpr half_t | ck::type_convert< half_t, bhalf_t > (bhalf_t x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::type_convert< bhalf_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ constexpr int8_t | ck::type_convert< int8_t, bhalf_t > (bhalf_t x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::type_convert< bhalf_t, int8_t > (int8_t x) |
| template<> | |
| __host__ __device__ constexpr f8_ocp_t | ck::type_convert< f8_ocp_t, int > (int x) |
| template<> | |
| __host__ __device__ constexpr bf8_ocp_t | ck::type_convert< bf8_ocp_t, int > (int x) |
| template<typename Y, enable_if_t< is_same_v< Y, ck::tf32_t >, bool > = false> | |
| __host__ __device__ constexpr float | ck::type_convert (float x) |
| template<typename Y, typename X> | |
| __host__ __device__ constexpr Y | ck::type_convert_sp (X x) |
| template<> | |
| __host__ __device__ constexpr int | ck::type_convert_sp< int, float > (float x) |
| template<> | |
| __host__ __device__ constexpr float | ck::type_convert_sp< float, int > (int x) |
| template<> | |
| __host__ __device__ constexpr int | ck::type_convert_sp< int, half_t > (half_t x) |
| template<> | |
| __host__ __device__ constexpr half_t | ck::type_convert_sp< half_t, int > (int x) |
| template<> | |
| __host__ __device__ constexpr int | ck::type_convert_sp< int, f8_t > (f8_t x) |
| template<> | |
| __host__ __device__ constexpr f8_t | ck::type_convert_sp< f8_t, int > (int x) |
| template<> | |
| __host__ __device__ constexpr int | ck::type_convert_sp< int, bhalf_t > (bhalf_t x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::type_convert_sp< bhalf_t, int > (int x) |
| template<> | |
| __host__ __device__ constexpr bhalf_t | ck::type_convert_sp< bhalf_t, float > (float x) |
| template<> | |
| __host__ __device__ constexpr half_t | ck::type_convert_sp< half_t, float > (float x) |
| template<typename Y, typename X> | |
| __host__ __device__ constexpr Y | ck::f8_convert_sr (X x) |
| template<> | |
| __host__ __device__ f8_fnuz_t | ck::f8_convert_sr< f8_fnuz_t, float > (float x) |
| template<> | |
| __host__ __device__ f8_fnuz_t | ck::f8_convert_sr< f8_fnuz_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ bf8_fnuz_t | ck::f8_convert_sr< bf8_fnuz_t, float > (float x) |
| template<> | |
| __host__ __device__ bf8_fnuz_t | ck::f8_convert_sr< bf8_fnuz_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::f8_convert_sr< f8_ocp_t, float > (float x) |
| Converts a float to a 8-bit float type (f8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ f8x2_ocp_t | ck::f8_convert_sr< f8x2_ocp_t, float2_t > (float2_t x) |
| Converts a vector of 2 floats to a vector of 2 8-bit float types (f8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::f8_convert_sr< bf8_ocp_t, float > (float x) |
| Converts a float to a 8-bit float type (bf8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf8x2_ocp_t | ck::f8_convert_sr< bf8x2_ocp_t, float2_t > (float2_t x) |
| Converts a vector of 2 floats to a vector of 2 8-bit float types (bf8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::f8_convert_sr< f8_ocp_t, half_t > (half_t x) |
| Converts a half_t to a 8-bit float type (f8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ f8x2_ocp_t | ck::f8_convert_sr< f8x2_ocp_t, half2_t > (half2_t x) |
| Converts a vector of 2 half_t to a vector of 2 8-bit float types (f8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::f8_convert_sr< bf8_ocp_t, half_t > (half_t x) |
| Converts a half_t to a 8-bit half_t type (bf8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf8x2_ocp_t | ck::f8_convert_sr< bf8x2_ocp_t, half2_t > (half2_t x) |
| Converts a vector of 2 half_t to a vector of 2 8-bit float types (bf8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::f8_convert_sr< f8_ocp_t, bhalf_t > (bhalf_t x) |
| Converts a bhalf_t to a 8-bit float type (f8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ f8x2_ocp_t | ck::f8_convert_sr< f8x2_ocp_t, bhalf2_t > (bhalf2_t x) |
| Converts a vector of 2 bhalf_t to a vector of 2 8-bit float types (f8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::f8_convert_sr< bf8_ocp_t, bhalf_t > (bhalf_t x) |
| Converts a bhalf_t to a 8-bit half_t type (bf8_ocp_t) using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf8x2_ocp_t | ck::f8_convert_sr< bf8x2_ocp_t, bhalf2_t > (bhalf2_t x) |
| Converts a vector of 2 bhalf_t to a vector of 2 8-bit float types (bf8_ocp_t) using stochastic rounding. | |
| template<typename Y, typename X> | |
| __host__ __device__ constexpr Y | ck::f8_convert_rne (X x) |
| template<> | |
| __host__ __device__ f8_fnuz_t | ck::f8_convert_rne< f8_fnuz_t, float > (float x) |
| template<> | |
| __host__ __device__ f8_fnuz_t | ck::f8_convert_rne< f8_fnuz_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ bf8_fnuz_t | ck::f8_convert_rne< bf8_fnuz_t, float > (float x) |
| template<> | |
| __host__ __device__ bf8_fnuz_t | ck::f8_convert_rne< bf8_fnuz_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::f8_convert_rne< f8_ocp_t, float > (float x) |
| Converts a float to a 8-bit float type (f8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ f8x2_ocp_t | ck::f8_convert_rne< f8x2_ocp_t, float2_t > (float2_t x) |
| Converts a vector of 2 floats to a vector of 2 8-bit float types (f8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::f8_convert_rne< bf8_ocp_t, float > (float x) |
| Converts a float to a 8-bit float type (bf8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ bf8x2_ocp_t | ck::f8_convert_rne< bf8x2_ocp_t, float2_t > (float2_t x) |
| Converts a vector of 2 floats to a vector of 2 8-bit float types (bf8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::f8_convert_rne< f8_ocp_t, half_t > (half_t x) |
| Converts a half_t to a 8-bit float type (f8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ f8x2_ocp_t | ck::f8_convert_rne< f8x2_ocp_t, half2_t > (half2_t x) |
| Converts a vector of 2 half_t to a vector of 2 8-bit float types (f8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::f8_convert_rne< bf8_ocp_t, half_t > (half_t x) |
| Converts a half_t to a 8-bit half_t type (bf8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ bf8x2_ocp_t | ck::f8_convert_rne< bf8x2_ocp_t, half2_t > (half2_t x) |
| Converts a vector of 2 half_t to a vector of 2 8-bit float types (bf8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::f8_convert_rne< f8_ocp_t, bhalf_t > (bhalf_t x) |
| Converts a bhalf_t to a 8-bit float type (f8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ f8x2_ocp_t | ck::f8_convert_rne< f8x2_ocp_t, bhalf2_t > (bhalf2_t x) |
| Converts a vector of 2 bhalf_t to a vector of 2 8-bit float types (f8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::f8_convert_rne< bf8_ocp_t, bhalf_t > (bhalf_t x) |
| Converts a bhalf_t to a 8-bit half_t type (bf8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ bf8x2_ocp_t | ck::f8_convert_rne< bf8x2_ocp_t, bhalf2_t > (bhalf2_t x) |
| Converts a vector of 2 bhalf_t to a vector of 2 8-bit float types (bf8_ocp_t) using rounding to nearest/even. | |
| template<> | |
| __host__ __device__ f8_fnuz_t | ck::type_convert< f8_fnuz_t, float > (float x) |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, f8_fnuz_t > (f8_fnuz_t x) |
| template<> | |
| __host__ __device__ float2_t | ck::type_convert< float2_t, f8x2_fnuz_t > (f8x2_fnuz_t x) |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, f8_ocp_t > (f8_ocp_t x) |
| Converts a f8_ocp_t value to a float value. | |
| template<> | |
| __host__ __device__ float2_t | ck::type_convert< float2_t, f8x2_ocp_t > (f8x2_ocp_t x) |
| Converts a vector of 2 f8_ocp_t values to a vector of 2 float values. | |
| template<> | |
| __host__ __device__ half_t | ck::type_convert< half_t, f8_ocp_t > (f8_ocp_t x) |
| Converts a f8_ocp_t value to a half_t value. | |
| template<> | |
| __host__ __device__ half2_t | ck::type_convert< half2_t, f8x2_ocp_t > (f8x2_ocp_t x) |
| Converts a vector of 2 f8_ocp_t values to a vector of 2 half_t values. | |
| template<> | |
| __host__ __device__ bhalf_t | ck::type_convert< bhalf_t, f8_ocp_t > (f8_ocp_t x) |
| Converts a f8_ocp_t value to a bhalf_t value. | |
| template<> | |
| __host__ __device__ bhalf2_t | ck::type_convert< bhalf2_t, f8x2_ocp_t > (f8x2_ocp_t x) |
| Converts a vector of 2 f8_ocp_t values to a vector of 2 bhalf_t values. | |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, bf8_ocp_t > (bf8_ocp_t x) |
| Converts a bf8_ocp_t value to a float value. | |
| template<> | |
| __host__ __device__ float2_t | ck::type_convert< float2_t, bf8x2_ocp_t > (bf8x2_ocp_t x) |
| Converts a vector of 2 bf8_ocp_t values to a vector of 2 float values. | |
| template<> | |
| __host__ __device__ half_t | ck::type_convert< half_t, bf8_ocp_t > (bf8_ocp_t x) |
| Converts a bf8_ocp_t value to a half_t value. | |
| template<> | |
| __host__ __device__ half2_t | ck::type_convert< half2_t, bf8x2_ocp_t > (bf8x2_ocp_t x) |
| Converts a vector of 2 bf8_ocp_t values to a vector of 2 half_t values. | |
| template<> | |
| __host__ __device__ bhalf_t | ck::type_convert< bhalf_t, bf8_ocp_t > (bf8_ocp_t x) |
| Converts a bf8_ocp_t value to a bhalf_t value. | |
| template<> | |
| __host__ __device__ bhalf2_t | ck::type_convert< bhalf2_t, bf8x2_ocp_t > (bf8x2_ocp_t x) |
| Converts a vector of 2 bf8_ocp_t values to a vector of 2 bhalf_t values. | |
| template<> | |
| __host__ __device__ float2_t | ck::type_convert< float2_t, pk_i4_t > (pk_i4_t x) |
| template<> | |
| __host__ __device__ half2_t | ck::type_convert< half2_t, pk_i4_t > (pk_i4_t x) |
| template<> | |
| __host__ __device__ bhalf2_t | ck::type_convert< bhalf2_t, pk_i4_t > (pk_i4_t x) |
| template<> | |
| __host__ __device__ half2_t | ck::type_convert< half2_t, float2_t > (float2_t x) |
| template<> | |
| __host__ __device__ f8_fnuz_t | ck::type_convert< f8_fnuz_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::type_convert< f8_ocp_t, half_t > (half_t x) |
| Converts a half_t value to a f8_ocp_t value with rounding determined by a flag. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::type_convert< bf8_ocp_t, half_t > (half_t x) |
| Converts a half_t value to a bf8_ocp_t value with rounding determined by a flag. | |
| template<> | |
| __host__ __device__ half_t | ck::type_convert< half_t, f8_fnuz_t > (f8_fnuz_t x) |
| template<> | |
| __host__ __device__ bf8_fnuz_t | ck::type_convert< bf8_fnuz_t, float > (float x) |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::type_convert< f8_ocp_t, float > (float x) |
| Converts a float value to a f8_ocp_t value with rounding determined by a flag. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::type_convert< bf8_ocp_t, float > (float x) |
| Converts a float value to a bf8_ocp_t value with rounding determined by a flag. | |
| template<> | |
| __host__ __device__ f8_ocp_t | ck::type_convert< f8_ocp_t, bhalf_t > (bhalf_t x) |
| Converts a bhalf_t value to a f8_ocp_t value with rounding determined by a flag. | |
| template<> | |
| __host__ __device__ bf8_ocp_t | ck::type_convert< bf8_ocp_t, bhalf_t > (bhalf_t x) |
| Converts a bhalf_t value to a bf8_ocp_t value with rounding determined by a flag. | |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, bf8_fnuz_t > (bf8_fnuz_t x) |
| template<> | |
| __host__ __device__ bf8_fnuz_t | ck::type_convert< bf8_fnuz_t, half_t > (half_t x) |
| template<> | |
| __host__ __device__ half_t | ck::type_convert< half_t, bf8_fnuz_t > (bf8_fnuz_t x) |
| __host__ __device__ f4_t | ck::f4_convert_rne (float x, float scale=1.0f) |
| __host__ __device__ f4x2_t | ck::f4_convert_rne (float2_t x, float scale=1.0f) |
| __host__ __device__ f4_t | ck::f4_convert_sr (float x, float scale=1.0f) |
| __host__ __device__ f4x2_t | ck::f4_convert_sr (float2_t x, float scale=1.0f) |
| template<> | |
| __host__ __device__ f4_t | ck::type_convert< f4_t, float > (float x) |
| template<> | |
| __host__ __device__ f4x2_t | ck::type_convert< f4x2_t, float2_t > (float2_t x) |
| template<> | |
| __host__ __device__ f4x2_pk_t | ck::type_convert< f4x2_pk_t, float2_t > (float2_t x) |
| template<> | |
| __host__ __device__ f4x32_t | ck::type_convert< f4x32_t, float32_t > (float32_t x) |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, f4_t > (f4_t x) |
| template<> | |
| __host__ __device__ float2_t | ck::type_convert< float2_t, f4x2_t > (f4x2_t x) |
| template<> | |
| __host__ __device__ float32_t | ck::type_convert< float32_t, f4x32_t > (f4x32_t x) |
| __host__ __device__ f6_t | ck::f6_convert_rne (float x, float scale=1.0f) |
| Converts a float to a 6-bit float type (f6_t) using round-to-nearest-even. | |
| __host__ __device__ f6x32_t | ck::f6_convert_rne (float32_t x, float scale=1.0f) |
| Converts a 32-element single-precision float array into a packed 6-bit representation. | |
| __host__ __device__ f6_t | ck::f6_convert_sr (float x, float scale=1.0f) |
| Converts a float to the 6-bit floating-point type (f6_t) using stochastic rounding. | |
| __host__ __device__ f6x32_t | ck::f6_convert_sr (float32_t x, float scale=1.0f) |
| Converts a 32-element single-precision float array into a packed 6-bit representation. | |
| template<> | |
| __host__ __device__ f6_t | ck::type_convert< f6_t, float > (float x) |
| Specializes the type conversion template for converting a float into the 6-bit float type (f6_t). | |
| template<> | |
| __host__ __device__ f6x32_t | ck::type_convert< f6x32_t, float32_t > (float32_t x) |
| Specializes the type conversion template for converting a vector of 32 floats into the vector of 32 6-bit float types (f6x32_t). | |
| template<> | |
| __host__ __device__ f6x32_pk_t | ck::type_convert< f6x32_pk_t, float32_t > (float32_t x) |
| template<> | |
| __host__ __device__ f6x16_t | ck::type_convert< f6x16_t, float16_t > (float16_t x) |
| template<> | |
| __host__ __device__ f6x16_pk_t | ck::type_convert< f6x16_pk_t, float16_t > (float16_t x) |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, f6_t > (f6_t x) |
| Specializes the type conversion template for converting the 6-bit float type (f6_t) to float. | |
| template<> | |
| __host__ __device__ float32_t | ck::type_convert< float32_t, f6x32_t > (f6x32_t x) |
| Specializes the type conversion template for converting the vector of 32 6-bit float types (f6x32_t) to vector of 32 floats. | |
| template<> | |
| __host__ __device__ float16_t | ck::type_convert< float16_t, f6x16_t > (f6x16_t x) |
| template<> | |
| __host__ __device__ float16_t | ck::type_convert< float16_t, f6x16_pk_t > (f6x16_pk_t x) |
| __host__ __device__ bf6_t | ck::bf6_convert_rne (float x, float scale=1.0f) |
| Converts a float to the 6-bit BF6 type using round-to-nearest-even. | |
| __host__ __device__ bf6x32_t | ck::bf6_convert_rne (float32_t x, float scale=1.0f) |
| Converts a vector of 32 floats to the vector of 32 6-bit BF6 types using round-to-nearest-even. | |
| __host__ __device__ bf6_t | ck::bf6_convert_sr (float x, float scale=1.0f) |
| Converts a float to the 6-bit BF6 type using stochastic rounding. | |
| __host__ __device__ bf6x32_t | ck::bf6_convert_sr (float32_t x, float scale=1.0f) |
| Converts a vector of 32 floats to the vector of 32 6-bit BF6 types using stochastic rounding. | |
| template<> | |
| __host__ __device__ bf6_t | ck::type_convert< bf6_t, float > (float x) |
| Specializes float-to-bf6_t conversion. | |
| template<> | |
| __host__ __device__ bf6x32_t | ck::type_convert< bf6x32_t, float32_t > (float32_t x) |
| Specializes vector of 32 float-to-bf6_t conversion. | |
| template<> | |
| __host__ __device__ bf6x32_pk_t | ck::type_convert< bf6x32_pk_t, float32_t > (float32_t x) |
| template<> | |
| __host__ __device__ bf6x16_t | ck::type_convert< bf6x16_t, float16_t > (float16_t x) |
| template<> | |
| __host__ __device__ bf6x16_pk_t | ck::type_convert< bf6x16_pk_t, float16_t > (float16_t x) |
| template<> | |
| __host__ __device__ float | ck::type_convert< float, bf6_t > (bf6_t x) |
| Specializes the type conversion template for converting a bf6_t value to float. | |
| template<> | |
| __host__ __device__ float32_t | ck::type_convert< float32_t, bf6x32_t > (bf6x32_t x) |
| Specializes the type conversion template for converting a vector of 32 bf6_t values to vector of 32 floats. | |
| template<> | |
| __host__ __device__ float16_t | ck::type_convert< float16_t, bf6x16_t > (bf6x16_t x) |
| template<> | |
| __host__ __device__ float16_t | ck::type_convert< float16_t, bf6x16_pk_t > (bf6x16_pk_t x) |
| template<typename Y, typename X, size_t NumElems> | |
| __host__ __device__ void | ck::array_convert (std::array< Y, NumElems > &y, const std::array< X, NumElems > &x) |
| template<typename Y, typename X, index_t NumElems> | |
| __host__ __device__ void | ck::array_convert (Array< Y, NumElems > &y, const Array< X, NumElems > &x) |