type_convert.hpp File Reference

type_convert.hpp File Reference#

Composable Kernel: type_convert.hpp File Reference
type_convert.hpp File Reference

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)