8 #if defined(__HIPCC_RTC__) || defined(CK_CODE_GEN_RTC)
15 __host__ __device__
static constexpr
int32_t Lowest() noexcept {
return -2147483647 - 1; }
17 __host__ __device__
static constexpr
int32_t Min() noexcept {
return -2147483647 - 1; }
19 __host__ __device__
static constexpr
int32_t Max() noexcept {
return 2147483647; }
21 __host__ __device__
static constexpr
int32_t Infinity() noexcept {
return 0; }
28 __host__ __device__
static constexpr
int16_t Lowest() noexcept {
return -32768; }
30 __host__ __device__
static constexpr
int16_t Min() noexcept {
return -32768; }
32 __host__ __device__
static constexpr
int16_t Max() noexcept {
return 32767; }
34 __host__ __device__
static constexpr
int16_t Infinity() noexcept {
return 0; }
40 struct NumericLimits<
int8_t>
42 __host__ __device__
static constexpr
int8_t Lowest() noexcept {
return -128; }
44 __host__ __device__
static constexpr
int8_t Min() noexcept {
return -128; }
46 __host__ __device__
static constexpr
int8_t Max() noexcept {
return 127; }
48 __host__ __device__
static constexpr
int8_t Infinity() noexcept {
return 0; }
50 __host__ __device__
static constexpr
int8_t QuietNaN() {
return 0; }
56 __host__ __device__
static constexpr
uint32_t Lowest() noexcept {
return 0; }
58 __host__ __device__
static constexpr
uint32_t Min() noexcept {
return 0; }
60 __host__ __device__
static constexpr
uint32_t Max() noexcept {
return 4294967295U; }
62 __host__ __device__
static constexpr
uint32_t Infinity() noexcept {
return 0; }
70 __host__ __device__
static constexpr
uint16_t Lowest() noexcept {
return 0; }
72 __host__ __device__
static constexpr
uint16_t Min() noexcept {
return 0; }
74 __host__ __device__
static constexpr
uint16_t Max() noexcept {
return 65535U; }
76 __host__ __device__
static constexpr
uint16_t Infinity() noexcept {
return 0; }
82 struct NumericLimits<float>
84 static constexpr
unsigned int binary_min = 0x00800000;
85 static constexpr
unsigned int binary_max = 0x7F7FFFFF;
86 static constexpr
unsigned int binary_lowest = 0xFF7FFFFF;
87 static constexpr
unsigned int binary_qnan = 0xFFC00001;
88 static constexpr
unsigned int binary_inf = 0x7F800000;
90 __host__ __device__
static constexpr
float Min() {
return bit_cast<float>(binary_min); }
92 __host__ __device__
static constexpr
float Max() {
return bit_cast<float>(binary_max); }
94 __host__ __device__
static constexpr
float Lowest() {
return bit_cast<float>(binary_lowest); }
96 __host__ __device__
static constexpr
float QuietNaN() {
return bit_cast<float>(binary_qnan); }
98 __host__ __device__
static constexpr
float Infinity() {
return bit_cast<float>(binary_inf); }
102 struct NumericLimits<
half_t>
104 static constexpr
unsigned short binary_min = 0x0400;
105 static constexpr
unsigned short binary_max = 0x7BFF;
106 static constexpr
unsigned short binary_lowest = 0xFBFF;
107 static constexpr
unsigned short binary_qnan = 0x7FFF;
109 __host__ __device__
static constexpr
half_t Min() {
return bit_cast<half_t>(binary_min); }
111 __host__ __device__
static constexpr
half_t Max() {
return bit_cast<half_t>(binary_max); }
113 __host__ __device__
static constexpr
half_t Lowest() {
return bit_cast<half_t>(binary_lowest); }
115 __host__ __device__
static constexpr
half_t QuietNaN() {
return bit_cast<half_t>(binary_qnan); }
118 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
120 struct NumericLimits<
int4_t>
122 __host__ __device__
static constexpr
int4_t Min() {
return int4_t(-8); }
124 __host__ __device__
static constexpr
int4_t Max() {
return int4_t(7); }
134 static constexpr
uint8_t binary_min = 0x08;
135 static constexpr
uint8_t binary_max = 0x7F;
136 static constexpr
uint8_t binary_lowest = 0xFF;
137 static constexpr
uint8_t binary_qnan = 0x80;
157 static constexpr
uint8_t binary_min = 0x04;
158 static constexpr
uint8_t binary_max = 0x7F;
159 static constexpr
uint8_t binary_lowest = 0xFF;
160 static constexpr
uint8_t binary_qnan = 0x80;
177 struct NumericLimits<f8_ocp_t>
179 static constexpr
uint8_t binary_min = 0x08;
180 static constexpr
uint8_t binary_max = 0x7E;
181 static constexpr
uint8_t binary_lowest = 0xFE;
182 static constexpr
uint8_t binary_qnan = 0x7F;
184 __host__ __device__
static constexpr f8_ocp_t
Min() {
return bit_cast<f8_ocp_t>(binary_min); }
186 __host__ __device__
static constexpr f8_ocp_t
Max() {
return bit_cast<f8_ocp_t>(binary_max); }
188 __host__ __device__
static constexpr f8_ocp_t
Lowest()
190 return bit_cast<f8_ocp_t>(binary_lowest);
193 __host__ __device__
static constexpr f8_ocp_t
QuietNaN()
195 return bit_cast<f8_ocp_t>(binary_qnan);
200 struct NumericLimits<bf8_ocp_t>
202 static constexpr
uint8_t binary_min = 0x04;
203 static constexpr
uint8_t binary_max = 0x7B;
204 static constexpr
uint8_t binary_lowest = 0xFB;
205 static constexpr
uint8_t binary_qnan = 0x7D;
207 __host__ __device__
static constexpr bf8_ocp_t
Min() {
return bit_cast<bf8_ocp_t>(binary_min); }
209 __host__ __device__
static constexpr bf8_ocp_t
Max() {
return bit_cast<bf8_ocp_t>(binary_max); }
211 __host__ __device__
static constexpr bf8_ocp_t
Lowest()
213 return bit_cast<bf8_ocp_t>(binary_lowest);
216 __host__ __device__
static constexpr bf8_ocp_t
QuietNaN()
218 return bit_cast<bf8_ocp_t>(binary_qnan);
223 struct NumericLimits<
f4_t>
225 static constexpr
uint8_t binary_min_normal = 0x2;
226 static constexpr
uint8_t binary_max_normal = 0x7;
227 static constexpr
uint8_t binary_lowest_normal = 0xF;
228 static constexpr
uint8_t binary_min_subnorm = 0x1;
229 static constexpr
uint8_t binary_max_subnorm = 0x1;
231 static constexpr
float data_max_normal_number = 6;
232 static constexpr
float data_min_subnormal_number = 0.5;
234 __host__ __device__
static constexpr
f4_t Min() {
return f4_t(binary_min_normal); }
235 __host__ __device__
static constexpr
f4_t Max() {
return f4_t(binary_max_normal); }
236 __host__ __device__
static constexpr
f4_t Lowest() {
return f4_t(binary_lowest_normal); }
237 __host__ __device__
static constexpr
f4_t MinSubnorm() {
return f4_t(binary_min_subnorm); }
238 __host__ __device__
static constexpr
f4_t MaxSubnorm() {
return f4_t(binary_max_subnorm); }
240 __host__ __device__
static constexpr
float DataMaxNorm() {
return data_max_normal_number; }
241 __host__ __device__
static constexpr
float DataMinSubnorm()
243 return data_min_subnormal_number;
248 struct NumericLimits<
f6_t>
250 static constexpr
uint8_t binary_min_normal = 0x08;
251 static constexpr
uint8_t binary_max_normal = 0x1F;
252 static constexpr
uint8_t binary_lowest_normal = 0x3F;
253 static constexpr
uint8_t binary_min_subnorm = 0x01;
254 static constexpr
uint8_t binary_max_subnorm = 0x07;
256 static constexpr
float data_max_normal_number = 7.5;
257 static constexpr
float data_min_subnormal_number = 0.125;
259 __host__ __device__
static constexpr
f6_t Min() {
return f6_t(binary_min_normal & 0b111111); }
260 __host__ __device__
static constexpr
f6_t Max() {
return f6_t(binary_max_normal & 0b111111); }
261 __host__ __device__
static constexpr
f6_t Lowest()
263 return f6_t(binary_lowest_normal & 0b111111);
265 __host__ __device__
static constexpr
f6_t MinSubnorm()
267 return f6_t(binary_min_subnorm & 0b111111);
269 __host__ __device__
static constexpr
f6_t MaxSubnorm()
271 return f6_t(binary_max_subnorm & 0b111111);
274 __host__ __device__
static constexpr
float DataMaxNorm() {
return data_max_normal_number; }
275 __host__ __device__
static constexpr
float DataMinSubnorm()
277 return data_min_subnormal_number;
282 struct NumericLimits<
bf6_t>
284 static constexpr
uint8_t binary_min_normal = 0x08;
285 static constexpr
uint8_t binary_max_normal = 0x1F;
286 static constexpr
uint8_t binary_lowest_normal = 0x3F;
287 static constexpr
uint8_t binary_min_subnorm = 0x01;
288 static constexpr
uint8_t binary_max_subnorm = 0x03;
290 static constexpr
float data_max_normal_number = 28;
291 static constexpr
float data_min_subnormal_number = 0.0625;
293 __host__ __device__
static constexpr
bf6_t Min() {
return bf6_t(binary_min_normal); }
294 __host__ __device__
static constexpr
bf6_t Max() {
return bf6_t(binary_max_normal); }
295 __host__ __device__
static constexpr
bf6_t Lowest() {
return bf6_t(binary_lowest_normal); }
296 __host__ __device__
static constexpr
bf6_t MinSubnorm() {
return bf6_t(binary_min_subnorm); }
297 __host__ __device__
static constexpr
bf6_t MaxSubnorm() {
return bf6_t(binary_max_subnorm); }
299 __host__ __device__
static constexpr
float DataMaxNorm() {
return data_max_normal_number; }
300 __host__ __device__
static constexpr
float DataMinSubnorm()
302 return data_min_subnormal_number;
307 template <
typename T>
312 __host__ __device__
static constexpr T
Lowest() {
return std::numeric_limits<T>::lowest(); }
315 return std::numeric_limits<T>::quiet_NaN();
317 __host__ __device__
static constexpr T
Infinity() {
return std::numeric_limits<T>::infinity(); }
323 static constexpr
unsigned short binary_min = 0x0400;
324 static constexpr
unsigned short binary_max = 0x7BFF;
325 static constexpr
unsigned short binary_lowest = 0xFBFF;
326 static constexpr
unsigned short binary_qnan = 0x7FFF;
328 __host__ __device__
static constexpr
half_t Min() {
return bit_cast<half_t>(binary_min); }
330 __host__ __device__
static constexpr
half_t Max() {
return bit_cast<half_t>(binary_max); }
332 __host__ __device__
static constexpr
half_t Lowest() {
return bit_cast<half_t>(binary_lowest); }
334 __host__ __device__
static constexpr
half_t QuietNaN() {
return bit_cast<half_t>(binary_qnan); }
337 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
339 struct NumericLimits<
int4_t>
341 __host__ __device__
static constexpr
int4_t Min() {
return int4_t(-8); }
343 __host__ __device__
static constexpr
int4_t Max() {
return int4_t(7); }
355 static constexpr
uint8_t binary_lowest = 0xFF;
378 static constexpr
uint8_t binary_lowest = 0xFF;
400 static constexpr
uint8_t binary_lowest = 0xFE;
403 __host__ __device__
static constexpr
f8_ocp_t Min() {
return bit_cast<f8_ocp_t>(binary_min); }
405 __host__ __device__
static constexpr
f8_ocp_t Max() {
return bit_cast<f8_ocp_t>(binary_max); }
409 return bit_cast<f8_ocp_t>(binary_lowest);
414 return bit_cast<f8_ocp_t>(binary_qnan);
423 static constexpr
uint8_t binary_lowest = 0xFB;
426 __host__ __device__
static constexpr
bf8_ocp_t Min() {
return bit_cast<bf8_ocp_t>(binary_min); }
428 __host__ __device__
static constexpr
bf8_ocp_t Max() {
return bit_cast<bf8_ocp_t>(binary_max); }
432 return bit_cast<bf8_ocp_t>(binary_lowest);
437 return bit_cast<bf8_ocp_t>(binary_qnan);
444 static constexpr
uint8_t binary_min_normal = 0x2;
445 static constexpr
uint8_t binary_max_normal = 0x7;
446 static constexpr
uint8_t binary_lowest_normal = 0xF;
447 static constexpr
uint8_t binary_min_subnorm = 0x1;
448 static constexpr
uint8_t binary_max_subnorm = 0x1;
450 static constexpr
float data_max_normal_number = 6;
451 static constexpr
float data_min_subnormal_number = 0.5;
453 __host__ __device__
static constexpr
f4_t Min() {
return f4_t(binary_min_normal); }
454 __host__ __device__
static constexpr
f4_t Max() {
return f4_t(binary_max_normal); }
455 __host__ __device__
static constexpr
f4_t Lowest() {
return f4_t(binary_lowest_normal); }
459 __host__ __device__
static constexpr
float DataMaxNorm() {
return data_max_normal_number; }
462 return data_min_subnormal_number;
469 static constexpr
uint8_t binary_min_normal = 0x08;
470 static constexpr
uint8_t binary_max_normal = 0x1F;
471 static constexpr
uint8_t binary_lowest_normal = 0x3F;
472 static constexpr
uint8_t binary_min_subnorm = 0x01;
473 static constexpr
uint8_t binary_max_subnorm = 0x07;
475 static constexpr
float data_max_normal_number = 7.5;
476 static constexpr
float data_min_subnormal_number = 0.125;
478 __host__ __device__
static constexpr
f6_t Min() {
return f6_t(binary_min_normal & 0b111111); }
479 __host__ __device__
static constexpr
f6_t Max() {
return f6_t(binary_max_normal & 0b111111); }
482 return f6_t(binary_lowest_normal & 0b111111);
486 return f6_t(binary_min_subnorm & 0b111111);
490 return f6_t(binary_max_subnorm & 0b111111);
493 __host__ __device__
static constexpr
float DataMaxNorm() {
return data_max_normal_number; }
496 return data_min_subnormal_number;
503 static constexpr
uint8_t binary_min_normal = 0x08;
504 static constexpr
uint8_t binary_max_normal = 0x1F;
505 static constexpr
uint8_t binary_lowest_normal = 0x3F;
506 static constexpr
uint8_t binary_min_subnorm = 0x01;
507 static constexpr
uint8_t binary_max_subnorm = 0x03;
509 static constexpr
float data_max_normal_number = 28;
510 static constexpr
float data_min_subnormal_number = 0.0625;
512 __host__ __device__
static constexpr
bf6_t Min() {
return bf6_t(binary_min_normal); }
513 __host__ __device__
static constexpr
bf6_t Max() {
return bf6_t(binary_max_normal); }
514 __host__ __device__
static constexpr
bf6_t Lowest() {
return bf6_t(binary_lowest_normal); }
518 __host__ __device__
static constexpr
float DataMaxNorm() {
return data_max_normal_number; }
521 return data_min_subnormal_number;
__host__ constexpr __device__ T max(T x)
Definition: math.hpp:84
__host__ constexpr __device__ T min(T x)
Definition: math.hpp:116
unsigned _BitInt(4) f4_t
Definition: data_type.hpp:32
_BitInt(6) f6_t
Definition: data_type.hpp:33
_Float16 half_t
Definition: data_type.hpp:30
_BitInt(4) int4_t
Definition: data_type.hpp:31
unsigned _BitInt(6) bf6_t
Definition: data_type.hpp:34
unsigned _BitInt(8) bf8_fnuz_t
Definition: amd_ck_fp8.hpp:37
_BitInt(8) f8_fnuz_t
Definition: amd_ck_fp8.hpp:36
signed short int16_t
Definition: stdint.h:122
unsigned short uint16_t
Definition: stdint.h:125
unsigned int uint32_t
Definition: stdint.h:126
signed int int32_t
Definition: stdint.h:123
unsigned char uint8_t
Definition: stdint.h:124
signed char int8_t
Definition: stdint.h:121
__host__ static constexpr __device__ float DataMaxNorm()
Definition: numeric_limits.hpp:518
__host__ static constexpr __device__ bf6_t MinSubnorm()
Definition: numeric_limits.hpp:515
__host__ static constexpr __device__ float DataMinSubnorm()
Definition: numeric_limits.hpp:519
__host__ static constexpr __device__ bf6_t MaxSubnorm()
Definition: numeric_limits.hpp:516
__host__ static constexpr __device__ bf6_t Max()
Definition: numeric_limits.hpp:513
__host__ static constexpr __device__ bf6_t Lowest()
Definition: numeric_limits.hpp:514
__host__ static constexpr __device__ bf6_t Min()
Definition: numeric_limits.hpp:512
__host__ static constexpr __device__ bf8_fnuz_t QuietNaN()
Definition: numeric_limits.hpp:392
__host__ static constexpr __device__ bf8_fnuz_t Min()
Definition: numeric_limits.hpp:386
__host__ static constexpr __device__ bf8_fnuz_t Lowest()
Definition: numeric_limits.hpp:390
__host__ static constexpr __device__ bf8_fnuz_t Max()
Definition: numeric_limits.hpp:388
__host__ static constexpr __device__ bf8_ocp_t Min()
Definition: numeric_limits.hpp:426
__host__ static constexpr __device__ bf8_ocp_t Lowest()
Definition: numeric_limits.hpp:430
__host__ static constexpr __device__ bf8_ocp_t Max()
Definition: numeric_limits.hpp:428
__host__ static constexpr __device__ bf8_ocp_t QuietNaN()
Definition: numeric_limits.hpp:435
__host__ static constexpr __device__ e8m0_bexp_t Binary_2()
Definition: numeric_limits.hpp:543
__host__ static constexpr __device__ e8m0_bexp_t Binary_142()
Definition: numeric_limits.hpp:549
__host__ static constexpr __device__ e8m0_bexp_t Max()
Definition: numeric_limits.hpp:540
__host__ static constexpr __device__ e8m0_bexp_t QuietNaN()
Definition: numeric_limits.hpp:541
__host__ static constexpr __device__ e8m0_bexp_t Binary_135()
Definition: numeric_limits.hpp:545
__host__ static constexpr __device__ e8m0_bexp_t Min()
Definition: numeric_limits.hpp:539
__host__ static constexpr __device__ e8m0_bexp_t Binary_1()
Definition: numeric_limits.hpp:542
__host__ static constexpr __device__ e8m0_bexp_t Binary_3()
Definition: numeric_limits.hpp:544
__host__ static constexpr __device__ float DataMinSubnorm()
Definition: numeric_limits.hpp:460
__host__ static constexpr __device__ f4_t Min()
Definition: numeric_limits.hpp:453
__host__ static constexpr __device__ f4_t Lowest()
Definition: numeric_limits.hpp:455
__host__ static constexpr __device__ float DataMaxNorm()
Definition: numeric_limits.hpp:459
__host__ static constexpr __device__ f4_t Max()
Definition: numeric_limits.hpp:454
__host__ static constexpr __device__ f4_t MaxSubnorm()
Definition: numeric_limits.hpp:457
__host__ static constexpr __device__ f4_t MinSubnorm()
Definition: numeric_limits.hpp:456
__host__ static constexpr __device__ float DataMaxNorm()
Definition: numeric_limits.hpp:493
__host__ static constexpr __device__ f6_t MinSubnorm()
Definition: numeric_limits.hpp:484
__host__ static constexpr __device__ f6_t MaxSubnorm()
Definition: numeric_limits.hpp:488
__host__ static constexpr __device__ f6_t Min()
Definition: numeric_limits.hpp:478
__host__ static constexpr __device__ f6_t Max()
Definition: numeric_limits.hpp:479
__host__ static constexpr __device__ f6_t Lowest()
Definition: numeric_limits.hpp:480
__host__ static constexpr __device__ float DataMinSubnorm()
Definition: numeric_limits.hpp:494
__host__ static constexpr __device__ f8_fnuz_t QuietNaN()
Definition: numeric_limits.hpp:369
__host__ static constexpr __device__ f8_fnuz_t Min()
Definition: numeric_limits.hpp:363
__host__ static constexpr __device__ f8_fnuz_t Max()
Definition: numeric_limits.hpp:365
__host__ static constexpr __device__ f8_fnuz_t Lowest()
Definition: numeric_limits.hpp:367
__host__ static constexpr __device__ f8_ocp_t Min()
Definition: numeric_limits.hpp:403
__host__ static constexpr __device__ f8_ocp_t Max()
Definition: numeric_limits.hpp:405
__host__ static constexpr __device__ f8_ocp_t QuietNaN()
Definition: numeric_limits.hpp:412
__host__ static constexpr __device__ f8_ocp_t Lowest()
Definition: numeric_limits.hpp:407
__host__ static constexpr __device__ half_t Max()
Definition: numeric_limits.hpp:330
__host__ static constexpr __device__ half_t Lowest()
Definition: numeric_limits.hpp:332
__host__ static constexpr __device__ half_t Min()
Definition: numeric_limits.hpp:328
__host__ static constexpr __device__ half_t QuietNaN()
Definition: numeric_limits.hpp:334
Definition: numeric_limits.hpp:309
__host__ static constexpr __device__ T Lowest()
Definition: numeric_limits.hpp:312
__host__ static constexpr __device__ T Infinity()
Definition: numeric_limits.hpp:317
__host__ static constexpr __device__ T QuietNaN()
Definition: numeric_limits.hpp:313
__host__ static constexpr __device__ T Min()
Definition: numeric_limits.hpp:310
__host__ static constexpr __device__ T Max()
Definition: numeric_limits.hpp:311
Definition: amd_ck_fp8.hpp:344
Unsigned representation of a conventional biased Float32 exponent.
Definition: e8m0.hpp:25
Definition: amd_ck_fp8.hpp:298