11 #if defined(__gfx950__)
12 #define CK_TILE_FP4_CVT_DEVICE 1
14 #define CK_TILE_FP4_CVT_DEVICE 0
17 #define TEST_convert_with_table 0
22 using fp32x2_t =
float __attribute__((ext_vector_type(2)));
23 using fp16x2_t = _Float16 __attribute__((ext_vector_type(2)));
38 template <
typename T,
typename = std::enable_if_t<std::is_
integral_v<T>>>
79 return (x1 << 4) | (x0 & 0b00001111);
82 #if TEST_convert_with_table
83 static constexpr
float e2m1_to_fp32_table[16] = {
84 0, 0.5, 1, 1.5, 2, 3, 4, 6, -0, -0.5, -1, -1.5, -2, -3, -4, -6};
85 static constexpr
fp16_t e2m1_to_fp16_table[16] = {
86 bit_cast<fp16_t>(
static_cast<uint16_t>(0x0000)),
87 bit_cast<fp16_t>(
static_cast<uint16_t>(0x3800)),
88 bit_cast<fp16_t>(
static_cast<uint16_t>(0x3C00)),
89 bit_cast<fp16_t>(
static_cast<uint16_t>(0x3E00)),
90 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4000)),
91 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4200)),
92 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4400)),
93 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4600)),
94 bit_cast<fp16_t>(
static_cast<uint16_t>(0x8000)),
95 bit_cast<fp16_t>(
static_cast<uint16_t>(0xB800)),
96 bit_cast<fp16_t>(
static_cast<uint16_t>(0xBC00)),
97 bit_cast<fp16_t>(
static_cast<uint16_t>(0xBE00)),
98 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC000)),
99 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC200)),
100 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC400)),
101 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC600))
114 static constexpr
int exp = 2;
115 static constexpr
int mant = 1;
116 static constexpr
int bias = 1;
153 static_assert(I < 2,
"Index is out of range.");
157 return
data & 0b00001111;
162 #if CK_TILE_FP4_CVT_DEVICE
165 template <
typename T>
168 if constexpr(std::is_same_v<T, fp32_t>)
169 return fp32x2_t(__builtin_amdgcn_cvt_scalef32_pk_f32_fp4(src, scale, 0))[0];
170 else if constexpr(std::is_same_v<T, fp32x2_t>)
171 return __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(src, scale, 0);
172 else if constexpr(std::is_same_v<T, fp16_t>)
173 return fp16x2_t(__builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 0))[0];
174 else if constexpr(std::is_same_v<T, fp16x2_t>)
175 return __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 0);
176 else if constexpr(std::is_same_v<T, bf16_t>)
177 return bf16x2_t(__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 0))[0];
178 else if constexpr(std::is_same_v<T, bf16x2_t>)
179 return __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 0);
184 template <
typename T>
192 if constexpr(std::is_same_v<T, fp32_t>)
193 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(cvt.u32, src, src, scale, 0);
194 else if constexpr(std::is_same_v<T, fp32x2_t>)
195 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(cvt.u32, src[0], src[1], scale, 0);
196 else if constexpr(std::is_same_v<T, fp16_t>)
197 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(cvt.u32,
fp16x2_t{src, src}, scale, 0);
198 else if constexpr(std::is_same_v<T, fp16x2_t>)
199 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(cvt.u32, src, scale, 0);
200 else if constexpr(std::is_same_v<T, bf16_t>)
201 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(cvt.u32,
bf16x2_t{src, src}, scale, 0);
202 else if constexpr(std::is_same_v<T, bf16x2_t>)
203 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(cvt.u32, src, scale, 0);
214 #if CK_TILE_FP4_CVT_DEVICE
215 return impl::_from_f4<bf16_t>(
data, scale);
223 #if CK_TILE_FP4_CVT_DEVICE
224 return impl::_from_f4<bf16x2_t>(
data, scale);
227 type_convert<bf16_t>(convert_to_float<pk_fp4_t>(
_unpack(
number<1>{}), scale))};
234 #if CK_TILE_FP4_CVT_DEVICE
235 return impl::_to_f4(x, scale);
237 return convert_to_type<pk_fp4_t>(x, scale);
242 #if CK_TILE_FP4_CVT_DEVICE
243 return impl::_to_f4(x, scale);
245 auto res = convert_to_type<pk_fp4_t>(x, scale);
251 #if CK_TILE_FP4_CVT_DEVICE
252 return impl::_to_f4(x, scale);
260 #if CK_TILE_FP4_CVT_DEVICE
261 return impl::_to_f4(x, scale);
269 #if CK_TILE_FP4_CVT_DEVICE
270 return impl::_to_f4(x, scale);
277 #if CK_TILE_FP4_CVT_DEVICE
278 return impl::_to_f4(x, scale);
285 #if CK_TILE_FP4_CVT_DEVICE
286 return impl::_to_f4(x, scale);
317 #if TEST_convert_with_table == 0
320 #if CK_TILE_FP4_CVT_DEVICE
321 return impl::_from_f4<fp32_t>(
data, scale);
328 #if CK_TILE_FP4_CVT_DEVICE
329 return impl::_from_f4<fp32x2_t>(
data, scale);
338 #if CK_TILE_FP4_CVT_DEVICE
339 return impl::_from_f4<fp16_t>(
data, scale);
346 #if CK_TILE_FP4_CVT_DEVICE
347 return impl::_from_f4<fp16x2_t>(
data, scale);
350 type_convert<fp16_t>(convert_to_float<pk_fp4_t>(
_unpack(
number<1>{}), scale))};
360 return fp32x2_t{e2m1_to_fp32_table[
_unpack(number<0>{})] * scale, e2m1_to_fp32_table[
_unpack(number<1>{}] * scale};
364 return type_convert<float>(e2m1_to_fp16_table[
_unpack(number<0>{})]) * scale;
369 type_convert<fp16_t>(type_convert<float>(e2m1_to_fp16_table[
_unpack(number<0>{})]) * scale),
370 type_convert<fp16_t>(type_convert<float>(e2m1_to_fp16_table[
_unpack(number<1>{})]) *
#define CK_TILE_DEVICE
Definition: config.hpp:41
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:42
Definition: cluster_descriptor.hpp:13
typename pk_fp4_t::type pk_fp4_raw_t
Definition: pk_fp4.hpp:107
ushort bfloat16_t
Definition: bfloat16.hpp:111
constexpr CK_TILE_HOST_DEVICE pk_fp4_t fp16_to_pk_fp4(const fp16_t &x, float scale)
Definition: pk_fp4.hpp:249
bfloat16_t bf16x2_t
Definition: pk_fp4.hpp:24
_Float16 fp16_t
Definition: half.hpp:110
float fp32x2_t
Definition: pk_fp4.hpp:22
bfloat16_t bf16_t
Definition: bfloat16.hpp:113
constexpr CK_TILE_HOST_DEVICE pk_fp4_t fp32x2_to_pk_fp4(const fp32x2_t &x, float scale)
Definition: pk_fp4.hpp:283
pk_float4_e2m1_t pk_fp4_t
Definition: pk_fp4.hpp:106
float fp32_t
Definition: pk_fp4.hpp:21
_Float16 fp16x2_t
Definition: half.hpp:385
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t float_to_pk_fp4(const float &x, float scale=1.f)
Definition: pk_fp4.hpp:240
constexpr CK_TILE_HOST_DEVICE float pk_fp4_to_float(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:304
constexpr CK_TILE_HOST_DEVICE fp16_t pk_fp4_to_fp16(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:308
CK_TILE_DEVICE bfloat16_t exp(bfloat16_t x)
Definition: bfloat16.hpp:411
constexpr CK_TILE_HOST_DEVICE pk_fp4_t bf16x2_to_pk_fp4(const bf16x2_t &x, float scale)
Definition: pk_fp4.hpp:275
constexpr CK_TILE_HOST_DEVICE pk_fp4_t fp16x2_to_pk_fp4(const fp16x2_t &x, float scale)
Definition: pk_fp4.hpp:267
constexpr CK_TILE_HOST_DEVICE fp32x2_t pk_fp4_to_fp32x2(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:292
constexpr CK_TILE_HOST_DEVICE pk_fp4_raw_t float_to_mxfp4(float x, float scale)
Definition: pk_fp4.hpp:232
constexpr CK_TILE_HOST_DEVICE pk_fp4_t bf16_to_pk_fp4(const bf16_t &x, float scale)
Definition: pk_fp4.hpp:258
constexpr CK_TILE_HOST_DEVICE fp16x2_t pk_fp4_to_fp16x2(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:296
constexpr CK_TILE_HOST_DEVICE bf16_t pk_fp4_to_bf16(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:312
constexpr CK_TILE_HOST_DEVICE bf16x2_t pk_fp4_to_bf16x2(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:300
const GenericPointer< typename T::ValueType > T2 value
Definition: pointer.h:1350
unsigned short uint16_t
Definition: stdint.h:125
unsigned int uint32_t
Definition: stdint.h:126
unsigned char uint8_t
Definition: stdint.h:124
Definition: integral_constant.hpp:13
static constexpr CK_TILE_HOST_DEVICE bool has_inf()
Definition: pk_fp4.hpp:141
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t min()
Definition: pk_fp4.hpp:133
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t denorm_min()
Definition: pk_fp4.hpp:139
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t infinity()
Definition: pk_fp4.hpp:143
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t round_error()
Definition: pk_fp4.hpp:137
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t epsilon()
Definition: pk_fp4.hpp:136
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t zero()
Definition: pk_fp4.hpp:138
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t quiet_NaN()
Definition: pk_fp4.hpp:145
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t signaling_NaN()
Definition: pk_fp4.hpp:147
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t lowest()
Definition: pk_fp4.hpp:135
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t max()
Definition: pk_fp4.hpp:134
pk_fp4_raw_t bitwise_type
Definition: pk_fp4.hpp:112
Definition: numeric.hpp:81
static constexpr int PackedSize
Definition: numeric.hpp:82
Definition: numeric.hpp:18
static constexpr CK_TILE_HOST_DEVICE T max()
Definition: numeric.hpp:26
Definition: pk_fp4.hpp:31
constexpr CK_TILE_HOST_DEVICE bf16x2_t to_bf16x2(float scale=1.f) const
Definition: pk_fp4.hpp:221
constexpr CK_TILE_HOST_DEVICE fp16x2_t to_fp16x2(float scale=1.f) const
Definition: pk_fp4.hpp:344
constexpr CK_TILE_HOST_DEVICE fp16_t to_fp16(float scale=1.f) const
Definition: pk_fp4.hpp:336
constexpr CK_TILE_HOST_DEVICE float to_float(float scale=1.f) const
Definition: pk_fp4.hpp:318
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t()
Definition: pk_fp4.hpp:37
uint8_t raw_type
Definition: pk_fp4.hpp:33
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t(float init, float scale=1.f)
Definition: pk_fp4.hpp:42
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t unpack(number< I >) const
Definition: pk_fp4.hpp:65
constexpr CK_TILE_HOST_DEVICE type & get()
Definition: pk_fp4.hpp:47
constexpr CK_TILE_HOST_DEVICE type _unpack(number< I >) const
constexpr CK_TILE_HOST_DEVICE type get() const
Definition: pk_fp4.hpp:48
constexpr CK_TILE_HOST_DEVICE fp32x2_t to_fp32x2(float scale=1.f) const
Definition: pk_fp4.hpp:326
constexpr CK_TILE_HOST_DEVICE bf16_t to_bf16(float scale=1.f) const
Definition: pk_fp4.hpp:212
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t(T init)
Definition: pk_fp4.hpp:39
type data
Definition: pk_fp4.hpp:35
raw_type type
Definition: pk_fp4.hpp:34
constexpr static CK_TILE_HOST_DEVICE pk_float4_e2m1_t pack(const pk_float4_e2m1_t &x0, const pk_float4_e2m1_t &x1)
Definition: pk_fp4.hpp:69
constexpr static CK_TILE_HOST_DEVICE type _pack(const type x0, const type x1)
Definition: pk_fp4.hpp:77
#define CK_TILE_ARITHMETIC_USING_FLOAT(attr_, type_)
Definition: numeric.hpp:106