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)));
37 template <
typename T,
typename = std::enable_if_t<std::is_
integral_v<T>>>
67 return (x1 << 4) | (x0 & 0b00001111);
70 #if TEST_convert_with_table
71 static constexpr
float e2m1_to_fp32_table[16] = {
72 0, 0.5, 1, 1.5, 2, 3, 4, 6, -0, -0.5, -1, -1.5, -2, -3, -4, -6};
73 static constexpr
fp16_t e2m1_to_fp16_table[16] = {
74 bit_cast<fp16_t>(
static_cast<uint16_t>(0x0000)),
75 bit_cast<fp16_t>(
static_cast<uint16_t>(0x3800)),
76 bit_cast<fp16_t>(
static_cast<uint16_t>(0x3C00)),
77 bit_cast<fp16_t>(
static_cast<uint16_t>(0x3E00)),
78 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4000)),
79 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4200)),
80 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4400)),
81 bit_cast<fp16_t>(
static_cast<uint16_t>(0x4600)),
82 bit_cast<fp16_t>(
static_cast<uint16_t>(0x8000)),
83 bit_cast<fp16_t>(
static_cast<uint16_t>(0xB800)),
84 bit_cast<fp16_t>(
static_cast<uint16_t>(0xBC00)),
85 bit_cast<fp16_t>(
static_cast<uint16_t>(0xBE00)),
86 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC000)),
87 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC200)),
88 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC400)),
89 bit_cast<fp16_t>(
static_cast<uint16_t>(0xC600))
102 static constexpr
int exp = 2;
103 static constexpr
int mant = 1;
104 static constexpr
int bias = 1;
141 static_assert(I < 2,
"Index is out of range.");
145 return
data & 0b00001111;
150 #if CK_TILE_FP4_CVT_DEVICE
153 template <
typename T>
156 if constexpr(std::is_same_v<T, fp32_t>)
157 return fp32x2_t(__builtin_amdgcn_cvt_scalef32_pk_f32_fp4(src, scale, 0))[0];
158 else if constexpr(std::is_same_v<T, fp32x2_t>)
159 return __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(src, scale, 0);
160 else if constexpr(std::is_same_v<T, fp16_t>)
161 return fp16x2_t(__builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 0))[0];
162 else if constexpr(std::is_same_v<T, fp16x2_t>)
163 return __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(src, scale, 0);
164 else if constexpr(std::is_same_v<T, bf16_t>)
165 return bf16x2_t(__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 0))[0];
166 else if constexpr(std::is_same_v<T, bf16x2_t>)
167 return __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(src, scale, 0);
172 template <
typename T>
180 if constexpr(std::is_same_v<T, fp32_t>)
181 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(cvt.u32, src, src, scale, 0);
182 else if constexpr(std::is_same_v<T, fp32x2_t>)
183 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(cvt.u32, src[0], src[1], scale, 0);
184 else if constexpr(std::is_same_v<T, fp16_t>)
185 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(cvt.u32,
fp16x2_t{src, src}, scale, 0);
186 else if constexpr(std::is_same_v<T, fp16x2_t>)
187 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(cvt.u32, src, scale, 0);
188 else if constexpr(std::is_same_v<T, bf16_t>)
189 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(cvt.u32,
bf16x2_t{src, src}, scale, 0);
190 else if constexpr(std::is_same_v<T, bf16x2_t>)
191 cvt.u32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(cvt.u32, src, scale, 0);
202 #if CK_TILE_FP4_CVT_DEVICE
203 return impl::_from_f4<bf16_t>(
data, scale);
211 #if CK_TILE_FP4_CVT_DEVICE
212 return impl::_from_f4<bf16x2_t>(
data, scale);
215 type_convert<bf16_t>(convert_to_float<pk_fp4_t>(
unpack(
number<1>{}), scale))};
222 #if CK_TILE_FP4_CVT_DEVICE
223 return impl::_to_f4(x, scale);
225 return convert_to_type<pk_fp4_t>(x, scale);
234 #if CK_TILE_FP4_CVT_DEVICE
235 return impl::_to_f4(x, scale);
242 #if CK_TILE_FP4_CVT_DEVICE
243 return impl::_to_f4(x, scale);
250 #if CK_TILE_FP4_CVT_DEVICE
251 return impl::_to_f4(x, scale);
258 #if CK_TILE_FP4_CVT_DEVICE
259 return impl::_to_f4(x, scale);
266 #if CK_TILE_FP4_CVT_DEVICE
267 return impl::_to_f4(x, scale);
298 #if TEST_convert_with_table == 0
301 #if CK_TILE_FP4_CVT_DEVICE
302 return impl::_from_f4<fp32_t>(
data, scale);
309 #if CK_TILE_FP4_CVT_DEVICE
310 return impl::_from_f4<fp32x2_t>(
data, scale);
319 #if CK_TILE_FP4_CVT_DEVICE
320 return impl::_from_f4<fp16_t>(
data, scale);
327 #if CK_TILE_FP4_CVT_DEVICE
328 return impl::_from_f4<fp16x2_t>(
data, scale);
331 type_convert<fp16_t>(convert_to_float<pk_fp4_t>(
unpack(
number<1>{}), scale))};
341 return fp32x2_t{e2m1_to_fp32_table[
unpack(number<0>{})] * scale, e2m1_to_fp32_table[
unpack(number<1>{}] * scale};
345 return type_convert<float>(e2m1_to_fp16_table[
unpack(number<0>{})]) * scale;
350 type_convert<fp16_t>(type_convert<float>(e2m1_to_fp16_table[
unpack(number<0>{})]) * scale),
351 type_convert<fp16_t>(type_convert<float>(e2m1_to_fp16_table[
unpack(number<1>{})]) * scale)};
#define CK_TILE_DEVICE
Definition: config.hpp:41
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:42
Definition: cluster_descriptor.hpp:13
_BitInt(8) fp8_t
Definition: float8.hpp:204
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:232
bfloat16_t bf16x2_t
Definition: pk_fp4.hpp:24
constexpr CK_TILE_HOST_DEVICE pk_fp4_t float_to_pk_fp4(const float &x, float scale)
Definition: pk_fp4.hpp:228
_Float16 fp16_t
Definition: half.hpp:110
constexpr CK_TILE_HOST_DEVICE uint8_t float_to_e2m1(float x, float scale=1.f)
Definition: pk_fp4.hpp:220
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:264
pk_float4_e2m1_t pk_fp4_t
Definition: pk_fp4.hpp:94
float fp32_t
Definition: pk_fp4.hpp:21
_Float16 fp16x2_t
Definition: half.hpp:385
constexpr CK_TILE_HOST_DEVICE float pk_fp4_to_float(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:285
constexpr CK_TILE_HOST_DEVICE fp16_t pk_fp4_to_fp16(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:289
CK_TILE_DEVICE bfloat16_t exp(bfloat16_t x)
Definition: bfloat16.hpp:423
constexpr CK_TILE_HOST_DEVICE pk_fp4_t bf16x2_to_pk_fp4(const bf16x2_t &x, float scale)
Definition: pk_fp4.hpp:256
constexpr CK_TILE_HOST_DEVICE pk_fp4_t fp16x2_to_pk_fp4(const fp16x2_t &x, float scale)
Definition: pk_fp4.hpp:248
constexpr CK_TILE_HOST_DEVICE fp32x2_t pk_fp4_to_fp32x2(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:273
constexpr CK_TILE_HOST_DEVICE pk_fp4_t bf16_to_pk_fp4(const bf16_t &x, float scale)
Definition: pk_fp4.hpp:240
constexpr CK_TILE_HOST_DEVICE fp16x2_t pk_fp4_to_fp16x2(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:277
typename pk_fp4_t::raw_type pk_fp4_raw_t
Definition: pk_fp4.hpp:95
constexpr CK_TILE_HOST_DEVICE bf16_t pk_fp4_to_bf16(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:293
constexpr CK_TILE_HOST_DEVICE bf16x2_t pk_fp4_to_bf16x2(const pk_fp4_t &x, float scale)
Definition: pk_fp4.hpp:281
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:129
static constexpr CK_TILE_HOST_DEVICE fp8_t denorm_min()
Definition: pk_fp4.hpp:127
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t min()
Definition: pk_fp4.hpp:121
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t infinity()
Definition: pk_fp4.hpp:131
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t round_error()
Definition: pk_fp4.hpp:125
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t epsilon()
Definition: pk_fp4.hpp:124
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t zero()
Definition: pk_fp4.hpp:126
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t quiet_NaN()
Definition: pk_fp4.hpp:133
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t signaling_NaN()
Definition: pk_fp4.hpp:135
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t lowest()
Definition: pk_fp4.hpp:123
static constexpr CK_TILE_HOST_DEVICE pk_fp4_t max()
Definition: pk_fp4.hpp:122
pk_fp4_raw_t bitwise_type
Definition: pk_fp4.hpp:100
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:30
constexpr CK_TILE_HOST_DEVICE bf16x2_t to_bf16x2(float scale=1.f) const
Definition: pk_fp4.hpp:209
constexpr CK_TILE_HOST_DEVICE raw_type & get()
Definition: pk_fp4.hpp:46
raw_type data
Definition: pk_fp4.hpp:34
constexpr CK_TILE_HOST_DEVICE fp16x2_t to_fp16x2(float scale=1.f) const
Definition: pk_fp4.hpp:325
constexpr CK_TILE_HOST_DEVICE fp16_t to_fp16(float scale=1.f) const
Definition: pk_fp4.hpp:317
constexpr CK_TILE_HOST_DEVICE float to_float(float scale=1.f) const
Definition: pk_fp4.hpp:299
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t()
Definition: pk_fp4.hpp:36
uint8_t raw_type
Definition: pk_fp4.hpp:32
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t(float init, float scale=1.f)
Definition: pk_fp4.hpp:41
constexpr CK_TILE_HOST_DEVICE raw_type unpack(number< I >) const
constexpr CK_TILE_HOST_DEVICE raw_type get() const
Definition: pk_fp4.hpp:47
constexpr CK_TILE_HOST_DEVICE fp32x2_t to_fp32x2(float scale=1.f) const
Definition: pk_fp4.hpp:307
constexpr CK_TILE_HOST_DEVICE bf16_t to_bf16(float scale=1.f) const
Definition: pk_fp4.hpp:200
constexpr CK_TILE_HOST_DEVICE pk_float4_e2m1_t(T init)
Definition: pk_fp4.hpp:38
constexpr static CK_TILE_HOST_DEVICE pk_float4_e2m1_t pack(const type x0, const type x1)
Definition: pk_fp4.hpp:65
raw_type type
Definition: pk_fp4.hpp:33
#define CK_TILE_ARITHMETIC_USING_FLOAT(attr_, type_)
Definition: numeric.hpp:106