/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/numeric_limits.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/numeric_limits.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/numeric_limits.hpp Source File
numeric_limits.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
3 #pragma once
5 
6 namespace ck {
7 
8 #if defined(__HIPCC_RTC__) || defined(CK_CODE_GEN_RTC)
9 template <typename T>
10 struct NumericLimits;
11 
12 template <>
13 struct NumericLimits<int32_t>
14 {
15  __host__ __device__ static constexpr int32_t Lowest() noexcept { return -2147483647 - 1; }
16 
17  __host__ __device__ static constexpr int32_t Min() noexcept { return -2147483647 - 1; }
18 
19  __host__ __device__ static constexpr int32_t Max() noexcept { return 2147483647; }
20 
21  __host__ __device__ static constexpr int32_t Infinity() noexcept { return 0; }
22 
23  __host__ __device__ static constexpr int32_t QuietNaN() { return 0; }
24 };
25 template <>
26 struct NumericLimits<int16_t>
27 {
28  __host__ __device__ static constexpr int16_t Lowest() noexcept { return -32768; }
29 
30  __host__ __device__ static constexpr int16_t Min() noexcept { return -32768; }
31 
32  __host__ __device__ static constexpr int16_t Max() noexcept { return 32767; }
33 
34  __host__ __device__ static constexpr int16_t Infinity() noexcept { return 0; }
35 
36  __host__ __device__ static constexpr int16_t QuietNaN() { return 0; }
37 };
38 
39 template <>
40 struct NumericLimits<int8_t>
41 {
42  __host__ __device__ static constexpr int8_t Lowest() noexcept { return -128; }
43 
44  __host__ __device__ static constexpr int8_t Min() noexcept { return -128; }
45 
46  __host__ __device__ static constexpr int8_t Max() noexcept { return 127; }
47 
48  __host__ __device__ static constexpr int8_t Infinity() noexcept { return 0; }
49 
50  __host__ __device__ static constexpr int8_t QuietNaN() { return 0; }
51 };
52 
53 template <>
54 struct NumericLimits<uint32_t>
55 {
56  __host__ __device__ static constexpr uint32_t Lowest() noexcept { return 0; }
57 
58  __host__ __device__ static constexpr uint32_t Min() noexcept { return 0; }
59 
60  __host__ __device__ static constexpr uint32_t Max() noexcept { return 4294967295U; }
61 
62  __host__ __device__ static constexpr uint32_t Infinity() noexcept { return 0; }
63 
64  __host__ __device__ static constexpr uint32_t QuietNaN() { return 0; }
65 };
66 
67 template <>
68 struct NumericLimits<uint16_t>
69 {
70  __host__ __device__ static constexpr uint16_t Lowest() noexcept { return 0; }
71 
72  __host__ __device__ static constexpr uint16_t Min() noexcept { return 0; }
73 
74  __host__ __device__ static constexpr uint16_t Max() noexcept { return 65535U; }
75 
76  __host__ __device__ static constexpr uint16_t Infinity() noexcept { return 0; }
77 
78  __host__ __device__ static constexpr uint16_t QuietNaN() { return 0; }
79 };
80 
81 template <>
82 struct NumericLimits<float>
83 {
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;
89 
90  __host__ __device__ static constexpr float Min() { return bit_cast<float>(binary_min); }
91 
92  __host__ __device__ static constexpr float Max() { return bit_cast<float>(binary_max); }
93 
94  __host__ __device__ static constexpr float Lowest() { return bit_cast<float>(binary_lowest); }
95 
96  __host__ __device__ static constexpr float QuietNaN() { return bit_cast<float>(binary_qnan); }
97 
98  __host__ __device__ static constexpr float Infinity() { return bit_cast<float>(binary_inf); }
99 };
100 
101 template <>
102 struct NumericLimits<half_t>
103 {
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;
108 
109  __host__ __device__ static constexpr half_t Min() { return bit_cast<half_t>(binary_min); }
110 
111  __host__ __device__ static constexpr half_t Max() { return bit_cast<half_t>(binary_max); }
112 
113  __host__ __device__ static constexpr half_t Lowest() { return bit_cast<half_t>(binary_lowest); }
114 
115  __host__ __device__ static constexpr half_t QuietNaN() { return bit_cast<half_t>(binary_qnan); }
116 };
117 
118 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
119 template <>
120 struct NumericLimits<int4_t>
121 {
122  __host__ __device__ static constexpr int4_t Min() { return int4_t(-8); }
123 
124  __host__ __device__ static constexpr int4_t Max() { return int4_t(7); }
125 
126  __host__ __device__ static constexpr int4_t Lowest() { return int4_t(-8); }
127 };
128 #endif // CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
129 
130 template <>
131 struct NumericLimits<f8_fnuz_t>
132 {
133  // negative zero nan mode with exp bias = 8
134  static constexpr uint8_t binary_min = 0x08; // 0b00001000
135  static constexpr uint8_t binary_max = 0x7F; // 0b01111111
136  static constexpr uint8_t binary_lowest = 0xFF; // 0b11111111
137  static constexpr uint8_t binary_qnan = 0x80; // 0b10000000
138  // ieee mode with exp bias = 7
139  // static constexpr uint8_t binary_min = 0x08; // 0b00001000
140  // static constexpr uint8_t binary_max = 0x77; // 0b01110111
141  // static constexpr uint8_t binary_lowest = 0xF7; // 0b11110111
142  // static constexpr uint8_t binary_qnan = 0x79; // any sign, exp=1111, mant!=0
143 
144  __host__ __device__ static constexpr f8_fnuz_t Min() { return f8_fnuz_t(binary_min); }
145 
146  __host__ __device__ static constexpr f8_fnuz_t Max() { return f8_fnuz_t(binary_max); }
147 
148  __host__ __device__ static constexpr f8_fnuz_t Lowest() { return f8_fnuz_t(binary_lowest); }
149 
150  __host__ __device__ static constexpr f8_fnuz_t QuietNaN() { return f8_fnuz_t(binary_qnan); }
151 };
152 
153 template <>
154 struct NumericLimits<bf8_fnuz_t>
155 {
156  // negative zero nan mode with exp bias = 16
157  static constexpr uint8_t binary_min = 0x04; // 0b00000100
158  static constexpr uint8_t binary_max = 0x7F; // 0b01111111
159  static constexpr uint8_t binary_lowest = 0xFF; // 0b11111111
160  static constexpr uint8_t binary_qnan = 0x80; // 0b10000000
161  // ieee mode with exp bias = 15
162  // static constexpr uint8_t binary_min = 0x04; // 0b00000100
163  // static constexpr uint8_t binary_max = 0x7B; // 0b01111011
164  // static constexpr uint8_t binary_lowest = 0xFB; // 0b11111011
165  // static constexpr uint8_t binary_qnan = 0x79; // any sign, exp=1111, mant!=
166 
167  __host__ __device__ static constexpr bf8_fnuz_t Min() { return bf8_fnuz_t(binary_min); }
168 
169  __host__ __device__ static constexpr bf8_fnuz_t Max() { return bf8_fnuz_t(binary_max); }
170 
171  __host__ __device__ static constexpr bf8_fnuz_t Lowest() { return bf8_fnuz_t(binary_lowest); }
172 
173  __host__ __device__ static constexpr bf8_fnuz_t QuietNaN() { return bf8_fnuz_t(binary_qnan); }
174 };
175 
176 template <>
177 struct NumericLimits<f8_ocp_t>
178 {
179  static constexpr uint8_t binary_min = 0x08; // 0b00001000 = 2^-6
180  static constexpr uint8_t binary_max = 0x7E; // 0b01111110 = 448
181  static constexpr uint8_t binary_lowest = 0xFE; // 0b11111110 = -448
182  static constexpr uint8_t binary_qnan = 0x7F; // 0b01111111
183 
184  __host__ __device__ static constexpr f8_ocp_t Min() { return bit_cast<f8_ocp_t>(binary_min); }
185 
186  __host__ __device__ static constexpr f8_ocp_t Max() { return bit_cast<f8_ocp_t>(binary_max); }
187 
188  __host__ __device__ static constexpr f8_ocp_t Lowest()
189  {
190  return bit_cast<f8_ocp_t>(binary_lowest);
191  }
192 
193  __host__ __device__ static constexpr f8_ocp_t QuietNaN()
194  {
195  return bit_cast<f8_ocp_t>(binary_qnan);
196  }
197 };
198 
199 template <>
200 struct NumericLimits<bf8_ocp_t>
201 {
202  static constexpr uint8_t binary_min = 0x04; // 0b00000100 = 2^-14
203  static constexpr uint8_t binary_max = 0x7B; // 0b01111011 = 57344
204  static constexpr uint8_t binary_lowest = 0xFB; // 0b11111011 = -57344
205  static constexpr uint8_t binary_qnan = 0x7D; // 0b01111101
206 
207  __host__ __device__ static constexpr bf8_ocp_t Min() { return bit_cast<bf8_ocp_t>(binary_min); }
208 
209  __host__ __device__ static constexpr bf8_ocp_t Max() { return bit_cast<bf8_ocp_t>(binary_max); }
210 
211  __host__ __device__ static constexpr bf8_ocp_t Lowest()
212  {
213  return bit_cast<bf8_ocp_t>(binary_lowest);
214  }
215 
216  __host__ __device__ static constexpr bf8_ocp_t QuietNaN()
217  {
218  return bit_cast<bf8_ocp_t>(binary_qnan);
219  }
220 };
221 
222 template <>
223 struct NumericLimits<f4_t>
224 {
225  static constexpr uint8_t binary_min_normal = 0x2; // 0b0010
226  static constexpr uint8_t binary_max_normal = 0x7; // 0b0111
227  static constexpr uint8_t binary_lowest_normal = 0xF; // 0b1111
228  static constexpr uint8_t binary_min_subnorm = 0x1; // 0b0001
229  static constexpr uint8_t binary_max_subnorm = 0x1; // 0b0001
230 
231  static constexpr float data_max_normal_number = 6;
232  static constexpr float data_min_subnormal_number = 0.5;
233 
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); }
239 
240  __host__ __device__ static constexpr float DataMaxNorm() { return data_max_normal_number; }
241  __host__ __device__ static constexpr float DataMinSubnorm()
242  {
243  return data_min_subnormal_number;
244  }
245 };
246 
247 template <>
248 struct NumericLimits<f6_t>
249 {
250  static constexpr uint8_t binary_min_normal = 0x08; // 0b001000
251  static constexpr uint8_t binary_max_normal = 0x1F; // 0b011111
252  static constexpr uint8_t binary_lowest_normal = 0x3F; // 0b111111
253  static constexpr uint8_t binary_min_subnorm = 0x01; // 0b000001
254  static constexpr uint8_t binary_max_subnorm = 0x07; // 0b000111
255 
256  static constexpr float data_max_normal_number = 7.5;
257  static constexpr float data_min_subnormal_number = 0.125;
258 
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()
262  {
263  return f6_t(binary_lowest_normal & 0b111111);
264  }
265  __host__ __device__ static constexpr f6_t MinSubnorm()
266  {
267  return f6_t(binary_min_subnorm & 0b111111);
268  }
269  __host__ __device__ static constexpr f6_t MaxSubnorm()
270  {
271  return f6_t(binary_max_subnorm & 0b111111);
272  }
273 
274  __host__ __device__ static constexpr float DataMaxNorm() { return data_max_normal_number; }
275  __host__ __device__ static constexpr float DataMinSubnorm()
276  {
277  return data_min_subnormal_number;
278  }
279 };
280 
281 template <>
282 struct NumericLimits<bf6_t>
283 {
284  static constexpr uint8_t binary_min_normal = 0x08; // 0b001000
285  static constexpr uint8_t binary_max_normal = 0x1F; // 0b011111
286  static constexpr uint8_t binary_lowest_normal = 0x3F; // 0b111111
287  static constexpr uint8_t binary_min_subnorm = 0x01; // 0b000001
288  static constexpr uint8_t binary_max_subnorm = 0x03; // 0b000011
289 
290  static constexpr float data_max_normal_number = 28;
291  static constexpr float data_min_subnormal_number = 0.0625;
292 
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); }
298 
299  __host__ __device__ static constexpr float DataMaxNorm() { return data_max_normal_number; }
300  __host__ __device__ static constexpr float DataMinSubnorm()
301  {
302  return data_min_subnormal_number;
303  }
304 };
305 
306 #else
307 template <typename T>
309 {
310  __host__ __device__ static constexpr T Min() { return std::numeric_limits<T>::min(); }
311  __host__ __device__ static constexpr T Max() { return std::numeric_limits<T>::max(); }
312  __host__ __device__ static constexpr T Lowest() { return std::numeric_limits<T>::lowest(); }
313  __host__ __device__ static constexpr T QuietNaN()
314  {
315  return std::numeric_limits<T>::quiet_NaN();
316  }
317  __host__ __device__ static constexpr T Infinity() { return std::numeric_limits<T>::infinity(); }
318 };
319 
320 template <>
322 {
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;
327 
328  __host__ __device__ static constexpr half_t Min() { return bit_cast<half_t>(binary_min); }
329 
330  __host__ __device__ static constexpr half_t Max() { return bit_cast<half_t>(binary_max); }
331 
332  __host__ __device__ static constexpr half_t Lowest() { return bit_cast<half_t>(binary_lowest); }
333 
334  __host__ __device__ static constexpr half_t QuietNaN() { return bit_cast<half_t>(binary_qnan); }
335 };
336 
337 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
338 template <>
339 struct NumericLimits<int4_t>
340 {
341  __host__ __device__ static constexpr int4_t Min() { return int4_t(-8); }
342 
343  __host__ __device__ static constexpr int4_t Max() { return int4_t(7); }
344 
345  __host__ __device__ static constexpr int4_t Lowest() { return int4_t(-8); }
346 };
347 #endif // CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
348 
349 template <>
351 {
352  // negative zero nan mode with exp bias = 8
353  static constexpr uint8_t binary_min = 0x08; // 0b00001000
354  static constexpr uint8_t binary_max = 0x7F; // 0b01111111
355  static constexpr uint8_t binary_lowest = 0xFF; // 0b11111111
356  static constexpr uint8_t binary_qnan = 0x80; // 0b10000000
357  // ieee mode with exp bias = 7
358  // static constexpr uint8_t binary_min = 0x08; // 0b00001000
359  // static constexpr uint8_t binary_max = 0x77; // 0b01110111
360  // static constexpr uint8_t binary_lowest = 0xF7; // 0b11110111
361  // static constexpr uint8_t binary_qnan = 0x79; // any sign, exp=1111, mant!=0
362 
363  __host__ __device__ static constexpr f8_fnuz_t Min() { return f8_fnuz_t(binary_min); }
364 
365  __host__ __device__ static constexpr f8_fnuz_t Max() { return f8_fnuz_t(binary_max); }
366 
367  __host__ __device__ static constexpr f8_fnuz_t Lowest() { return f8_fnuz_t(binary_lowest); }
368 
369  __host__ __device__ static constexpr f8_fnuz_t QuietNaN() { return f8_fnuz_t(binary_qnan); }
370 };
371 
372 template <>
374 {
375  // negative zero nan mode with exp bias = 16
376  static constexpr uint8_t binary_min = 0x04; // 0b00000100
377  static constexpr uint8_t binary_max = 0x7F; // 0b01111111
378  static constexpr uint8_t binary_lowest = 0xFF; // 0b11111111
379  static constexpr uint8_t binary_qnan = 0x80; // 0b10000000
380  // ieee mode with exp bias = 15
381  // static constexpr uint8_t binary_min = 0x04; // 0b00000100
382  // static constexpr uint8_t binary_max = 0x7B; // 0b01111011
383  // static constexpr uint8_t binary_lowest = 0xFB; // 0b11111011
384  // static constexpr uint8_t binary_qnan = 0x79; // any sign, exp=1111, mant!=
385 
386  __host__ __device__ static constexpr bf8_fnuz_t Min() { return bf8_fnuz_t(binary_min); }
387 
388  __host__ __device__ static constexpr bf8_fnuz_t Max() { return bf8_fnuz_t(binary_max); }
389 
390  __host__ __device__ static constexpr bf8_fnuz_t Lowest() { return bf8_fnuz_t(binary_lowest); }
391 
392  __host__ __device__ static constexpr bf8_fnuz_t QuietNaN() { return bf8_fnuz_t(binary_qnan); }
393 };
394 
395 template <>
397 {
398  static constexpr uint8_t binary_min = 0x08; // 0b00001000 = 2^-6
399  static constexpr uint8_t binary_max = 0x7E; // 0b01111110 = 448
400  static constexpr uint8_t binary_lowest = 0xFE; // 0b11111110 = -448
401  static constexpr uint8_t binary_qnan = 0x7F; // 0b01111111
402 
403  __host__ __device__ static constexpr f8_ocp_t Min() { return bit_cast<f8_ocp_t>(binary_min); }
404 
405  __host__ __device__ static constexpr f8_ocp_t Max() { return bit_cast<f8_ocp_t>(binary_max); }
406 
407  __host__ __device__ static constexpr f8_ocp_t Lowest()
408  {
409  return bit_cast<f8_ocp_t>(binary_lowest);
410  }
411 
412  __host__ __device__ static constexpr f8_ocp_t QuietNaN()
413  {
414  return bit_cast<f8_ocp_t>(binary_qnan);
415  }
416 };
417 
418 template <>
420 {
421  static constexpr uint8_t binary_min = 0x04; // 0b00000100 = 2^-14
422  static constexpr uint8_t binary_max = 0x7B; // 0b01111011 = 57344
423  static constexpr uint8_t binary_lowest = 0xFB; // 0b11111011 = -57344
424  static constexpr uint8_t binary_qnan = 0x7D; // 0b01111101
425 
426  __host__ __device__ static constexpr bf8_ocp_t Min() { return bit_cast<bf8_ocp_t>(binary_min); }
427 
428  __host__ __device__ static constexpr bf8_ocp_t Max() { return bit_cast<bf8_ocp_t>(binary_max); }
429 
430  __host__ __device__ static constexpr bf8_ocp_t Lowest()
431  {
432  return bit_cast<bf8_ocp_t>(binary_lowest);
433  }
434 
435  __host__ __device__ static constexpr bf8_ocp_t QuietNaN()
436  {
437  return bit_cast<bf8_ocp_t>(binary_qnan);
438  }
439 };
440 
441 template <>
443 {
444  static constexpr uint8_t binary_min_normal = 0x2; // 0b0010
445  static constexpr uint8_t binary_max_normal = 0x7; // 0b0111
446  static constexpr uint8_t binary_lowest_normal = 0xF; // 0b1111
447  static constexpr uint8_t binary_min_subnorm = 0x1; // 0b0001
448  static constexpr uint8_t binary_max_subnorm = 0x1; // 0b0001
449 
450  static constexpr float data_max_normal_number = 6;
451  static constexpr float data_min_subnormal_number = 0.5;
452 
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); }
456  __host__ __device__ static constexpr f4_t MinSubnorm() { return f4_t(binary_min_subnorm); }
457  __host__ __device__ static constexpr f4_t MaxSubnorm() { return f4_t(binary_max_subnorm); }
458 
459  __host__ __device__ static constexpr float DataMaxNorm() { return data_max_normal_number; }
460  __host__ __device__ static constexpr float DataMinSubnorm()
461  {
462  return data_min_subnormal_number;
463  }
464 };
465 
466 template <>
468 {
469  static constexpr uint8_t binary_min_normal = 0x08; // 0b001000
470  static constexpr uint8_t binary_max_normal = 0x1F; // 0b011111
471  static constexpr uint8_t binary_lowest_normal = 0x3F; // 0b111111
472  static constexpr uint8_t binary_min_subnorm = 0x01; // 0b000001
473  static constexpr uint8_t binary_max_subnorm = 0x07; // 0b000111
474 
475  static constexpr float data_max_normal_number = 7.5;
476  static constexpr float data_min_subnormal_number = 0.125;
477 
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); }
480  __host__ __device__ static constexpr f6_t Lowest()
481  {
482  return f6_t(binary_lowest_normal & 0b111111);
483  }
484  __host__ __device__ static constexpr f6_t MinSubnorm()
485  {
486  return f6_t(binary_min_subnorm & 0b111111);
487  }
488  __host__ __device__ static constexpr f6_t MaxSubnorm()
489  {
490  return f6_t(binary_max_subnorm & 0b111111);
491  }
492 
493  __host__ __device__ static constexpr float DataMaxNorm() { return data_max_normal_number; }
494  __host__ __device__ static constexpr float DataMinSubnorm()
495  {
496  return data_min_subnormal_number;
497  }
498 };
499 
500 template <>
502 {
503  static constexpr uint8_t binary_min_normal = 0x08; // 0b001000
504  static constexpr uint8_t binary_max_normal = 0x1F; // 0b011111
505  static constexpr uint8_t binary_lowest_normal = 0x3F; // 0b111111
506  static constexpr uint8_t binary_min_subnorm = 0x01; // 0b000001
507  static constexpr uint8_t binary_max_subnorm = 0x03; // 0b000011
508 
509  static constexpr float data_max_normal_number = 28;
510  static constexpr float data_min_subnormal_number = 0.0625;
511 
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); }
515  __host__ __device__ static constexpr bf6_t MinSubnorm() { return bf6_t(binary_min_subnorm); }
516  __host__ __device__ static constexpr bf6_t MaxSubnorm() { return bf6_t(binary_max_subnorm); }
517 
518  __host__ __device__ static constexpr float DataMaxNorm() { return data_max_normal_number; }
519  __host__ __device__ static constexpr float DataMinSubnorm()
520  {
521  return data_min_subnormal_number;
522  }
523 };
524 
525 #endif
526 
527 template <>
529 {
530  static constexpr e8m0_bexp_t binary_min = 0x00; // 0b00000000
531  static constexpr e8m0_bexp_t binary_max = 0xFE; // 0b11111110
532  static constexpr e8m0_bexp_t binary_qnan = 0xFF; // 0b11111111
533  static constexpr e8m0_bexp_t binary_1 = 0x7F; // 0b01111111
534  static constexpr e8m0_bexp_t binary_2 = 0x80; // 0b10000000
535  static constexpr e8m0_bexp_t binary_3 = 0x82; // 0b10000010
536  static constexpr e8m0_bexp_t binary_135 = 0x87; // 0b10000111
537  static constexpr e8m0_bexp_t binary_142 = 0x8E; // 0b10001110
538 
539  __host__ __device__ static constexpr e8m0_bexp_t Min() { return e8m0_bexp_t(binary_min); }
540  __host__ __device__ static constexpr e8m0_bexp_t Max() { return e8m0_bexp_t(binary_max); }
541  __host__ __device__ static constexpr e8m0_bexp_t QuietNaN() { return e8m0_bexp_t(binary_qnan); }
542  __host__ __device__ static constexpr e8m0_bexp_t Binary_1() { return e8m0_bexp_t(binary_1); }
543  __host__ __device__ static constexpr e8m0_bexp_t Binary_2() { return e8m0_bexp_t(binary_2); }
544  __host__ __device__ static constexpr e8m0_bexp_t Binary_3() { return e8m0_bexp_t(binary_3); }
545  __host__ __device__ static constexpr e8m0_bexp_t Binary_135()
546  {
547  return e8m0_bexp_t(binary_135);
548  }
549  __host__ __device__ static constexpr e8m0_bexp_t Binary_142()
550  {
551  return e8m0_bexp_t(binary_142);
552  }
553 };
554 
555 } // namespace ck
__host__ constexpr __device__ T max(T x)
Definition: math.hpp:84
__host__ constexpr __device__ T min(T x)
Definition: math.hpp:116
Definition: ck.hpp:267
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