HDK
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
float8.h
Go to the documentation of this file.
1 // Copyright (c) Microsoft Corporation. All rights reserved.
2 // Licensed under the MIT License.
3 
4 #pragma once
5 
6 #if !defined(DISABLE_FLOAT8_TYPES)
7 
8 #include "endian.h"
9 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
10 #include "cuda_fp8.h"
11 #endif
12 
13 #if !defined(__CUDACC__) && !defined(__HIPCC__)
14 #include "core/common/narrow.h"
15 #endif
16 
17 #include "core/common/common.h"
18 
19 namespace onnxruntime {
20 
21 #if defined(__CUDACC__) || defined(__HIPCC__)
22 #define ORT_HOST_DEVICE __host__ __device__
23 #else
24 #define ORT_HOST_DEVICE
25 #endif
26 
27 // Float8E4M3FN
28 struct Float8E4M3FN {
29  uint8_t val{0};
30 #if defined(__HIP__)
31  ORT_HOST_DEVICE Float8E4M3FN() = default;
32 #else
33  Float8E4M3FN() = default;
34 #endif
35  struct FromBitsT {};
36  static constexpr ORT_HOST_DEVICE FromBitsT FromBits() { return FromBitsT(); }
37  constexpr ORT_HOST_DEVICE Float8E4M3FN(unsigned char bits, FromBitsT) : val(bits) {}
38 
39  inline explicit ORT_HOST_DEVICE Float8E4M3FN(float v, bool saturate = true) {
40 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
41  val = __nv_cvt_float_to_fp8(v, saturate ? __NV_SATFINITE : __NV_NOSAT, __NV_E4M3);
42 #else
43  uint32_t b;
44  std::memcpy(&b, &v, sizeof(b));
45 
46  val = static_cast<uint8_t>((b & 0x80000000) >> 24); // sign
47  if ((b & 0x7fffffff) == 0x7f800000) { // infinity
48  if (saturate) {
49  val |= 126;
50  } else {
51  val |= 0x7f;
52  }
53  } else if ((b & 0x7F800000) == 0x7F800000) { // NaN
54  val |= 0x7f;
55  } else {
56  uint8_t e = static_cast<uint8_t>((b & 0x7F800000) >> 23); // exponent
57  uint32_t m = static_cast<uint32_t>(b & 0x007FFFFF); // mantissa
58  if (e != 0) {
59  if (e < 117) {
60  } else if (e < 121) {
61  // denormalized number
62  auto d = 120 - e;
63  if (d < 3) {
64  val |= 1 << (2 - d);
65  val |= m >> (21 + d);
66  } else if (m > 0) {
67  val |= 1;
68  }
69  auto mask = 1 << (20 + d);
70  if ((m & mask) && ((val & 1) || ((m & (mask - 1)) > 0) || ((m & mask) && (m & (mask << 1)) && ((m & (mask - 1)) == 0)))) {
71  // rounding
72  val += 1;
73  }
74  } else if (e < 136) {
75  // normalized number
76  auto ex = e - 120;
77  if (ex == 0) {
78  val |= 0x4;
79  val |= m >> 21;
80  } else {
81  val |= ex << 3;
82  val |= m >> 20;
83  if ((val & 0x7F) == 0x7F) {
84  val &= 0xFE;
85  }
86  }
87  if ((m & 0x80000) && ((m & 0x100000) || (m & 0x7FFFF))) {
88  if ((val & 0x7F) < 0x7E) {
89  // rounding
90  val += 1;
91  } else if (!saturate) {
92  val |= 0x7F;
93  }
94  }
95  } else if (saturate) {
96  val |= 126; // 0b01111110
97  } else {
98  val |= 0x7F;
99  }
100  }
101  }
102 #endif
103  }
104 
105  inline ORT_HOST_DEVICE float ToFloat() const {
106 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
107  return __half2float(__nv_cvt_fp8_to_halfraw(val, __NV_E4M3));
108 #else
109  uint32_t res;
110  if (val == 255) {
111  res = 0xffc00000;
112  } else if (val == 127) {
113  res = 0x7fc00000;
114  } else {
115  uint32_t expo = (val & 0x78) >> 3;
116  uint32_t mant = val & 0x07;
117  uint32_t sign = val & 0x80;
118  res = sign << 24;
119  if (expo == 0) {
120  if (mant > 0) {
121  expo = 0x7F - 7;
122  if ((mant & 0x4) == 0) {
123  mant &= 0x3;
124  mant <<= 1;
125  expo -= 1;
126  }
127  if ((mant & 0x4) == 0) {
128  mant &= 0x3;
129  mant <<= 1;
130  expo -= 1;
131  }
132  res |= (mant & 0x3) << 21;
133  res |= expo << 23;
134  }
135  } else {
136  res |= mant << 20;
137  expo -= 0x7;
138  expo += 0x7F;
139  res |= expo << 23;
140  }
141  }
142  float float_res;
143  std::memcpy(&float_res, &res, sizeof(float));
144  return float_res;
145 #endif
146  }
147 
148  inline ORT_HOST_DEVICE operator float() const { return ToFloat(); }
149 
150 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
151  explicit ORT_HOST_DEVICE Float8E4M3FN(const __nv_fp8_e4m3& value) { val = *reinterpret_cast<const unsigned char*>(&value); }
152  explicit ORT_HOST_DEVICE operator __nv_fp8_e4m3() const { return *reinterpret_cast<const __nv_fp8_e4m3*>(&val); }
153 #endif
154 };
155 
156 inline ORT_HOST_DEVICE bool operator==(const Float8E4M3FN& left, const Float8E4M3FN& right) { return left.val == right.val; }
157 inline ORT_HOST_DEVICE bool operator!=(const Float8E4M3FN& left, const Float8E4M3FN& right) { return left.val != right.val; }
158 inline ORT_HOST_DEVICE bool operator<(const Float8E4M3FN& left, const Float8E4M3FN& right) { return left.val < right.val; }
159 
160 // User defined suffixes to make it easier to declare
161 // initializers with MLFloat8E4M3FN and Float8E4M3FN from unsigned char
162 #if !defined(__CUDACC__) && !defined(__HIPCC__)
163 
164 inline Float8E4M3FN operator"" _f8e4m3fn(unsigned long long int v) {
165  return Float8E4M3FN(narrow<uint8_t>(v), Float8E4M3FN::FromBits());
166 }
167 
168 inline Float8E4M3FN operator"" _f8e4m3fnp8(long double v) {
169  return Float8E4M3FN(static_cast<float>(v), true);
170 }
171 
172 #endif
173 
174 inline void Float8E4M3FNToFloat(const Float8E4M3FN* blf, float* flt, size_t size) {
175  auto src = blf;
176  auto d = flt;
177  for (; size != 0; ++src, ++d, --size) {
178  *d = src->ToFloat();
179  }
180 }
181 
182 inline void FloatToFloat8E4M3FN(const float* flt, Float8E4M3FN* blf, size_t size, bool saturate) {
183  auto src = flt;
184  auto d = blf;
185  for (; size != 0; ++src, ++d, --size) {
186  new (d) Float8E4M3FN(*src, saturate);
187  }
188 }
189 
190 // Float8E4M3FNUZ
192  uint8_t val{0};
193 #if defined(__HIP__)
194  ORT_HOST_DEVICE Float8E4M3FNUZ() = default;
195 #else
196  Float8E4M3FNUZ() = default;
197 #endif
198 
199  struct FromBitsT {};
200  static constexpr ORT_HOST_DEVICE FromBitsT FromBits() { return FromBitsT(); }
201  constexpr ORT_HOST_DEVICE Float8E4M3FNUZ(unsigned char bits, FromBitsT) : val(bits) {}
202 
203  inline explicit ORT_HOST_DEVICE Float8E4M3FNUZ(float v, bool saturate = true) {
204  // This type does not exist on CUDA.
205  uint32_t b;
206  std::memcpy(&b, &v, sizeof(b));
207 
208  val = static_cast<uint8_t>((b & 0x80000000) >> 24); // sign
209  if ((b & 0x7fffffff) == 0x7f800000) { // infinity
210  if (saturate) {
211  // the highest available value
212  val |= 0x7F;
213  } else {
214  // NaN
215  val = 0x80;
216  }
217  } else if ((b & 0x7F800000) == 0x7F800000) { // NaN
218  val = 0x80;
219  } else {
220  uint8_t e = static_cast<uint8_t>((b & 0x7F800000) >> 23); // exponent
221  uint32_t m = static_cast<uint32_t>(b & 0x007FFFFF); // mantissa
222  if (e != 0) {
223  if (e < 116) {
224  } else if (e < 120) {
225  // denormalized number
226  auto d = 119 - e;
227  if (d < 3) {
228  val |= 1 << (2 - d);
229  val |= m >> (21 + d);
230  } else if (m > 0) {
231  val |= 1;
232  }
233  auto mask = 1 << (20 + d);
234  if ((m & mask) && ((val & 1) || ((m & (mask - 1)) > 0) || ((m & mask) && (m & (mask << 1)) && ((m & (mask - 1)) == 0)))) {
235  // rounding
236  val += 1;
237  }
238  } else if (e < 135) {
239  // normalized number
240  auto ex = e - 119;
241  if (ex == 0) {
242  val |= 0x4;
243  val |= m >> 21;
244  } else {
245  val |= ex << 3;
246  val |= m >> 20;
247  }
248  if ((m & 0x80000) && ((m & 0x100000) || (m & 0x7FFFF))) {
249  if ((val & 0x7F) < 0x7F) {
250  // rounding
251  val += 1;
252  } else if (!saturate) {
253  val = 0x80;
254  }
255  }
256  } else if (saturate) {
257  val |= 0x7F;
258  } else {
259  val = 0x80;
260  }
261  } else if (m == 0) {
262  // -0
263  val = 0;
264  }
265  }
266  }
267 
268  inline ORT_HOST_DEVICE float ToFloat() const {
269  // This type does not exist on CUDA.
270  uint32_t res;
271  if (val == 0x80) {
272  res = 0xffc00000;
273  } else {
274  uint32_t expo = (val & 0x78) >> 3;
275  uint32_t mant = val & 0x07;
276  uint32_t sign = val & 0x80;
277  res = sign << 24;
278  if (expo == 0) {
279  if (mant > 0) {
280  expo = 0x7F - 8;
281  if ((mant & 0x4) == 0) {
282  mant &= 0x3;
283  mant <<= 1;
284  expo -= 1;
285  }
286  if ((mant & 0x4) == 0) {
287  mant &= 0x3;
288  mant <<= 1;
289  expo -= 1;
290  }
291  res |= (mant & 0x3) << 21;
292  res |= expo << 23;
293  }
294  } else {
295  res |= mant << 20;
296  expo -= 8;
297  expo += 0x7F;
298  res |= expo << 23;
299  }
300  }
301  float float_res;
302  std::memcpy(&float_res, &res, sizeof(float));
303  return float_res;
304  }
305 
306  inline ORT_HOST_DEVICE operator float() const { return ToFloat(); }
307 };
308 
309 inline ORT_HOST_DEVICE bool operator==(const Float8E4M3FNUZ& left, const Float8E4M3FNUZ& right) { return left.val == right.val; }
310 inline ORT_HOST_DEVICE bool operator!=(const Float8E4M3FNUZ& left, const Float8E4M3FNUZ& right) { return left.val != right.val; }
311 inline ORT_HOST_DEVICE bool operator<(const Float8E4M3FNUZ& left, const Float8E4M3FNUZ& right) { return left.val < right.val; }
312 
313 // User defined suffixes to make it easier to declare
314 // initializers with MLFloat8E4M3FN and Float8E4M3FN from unsigned char
315 #if !defined(__CUDACC__) && !defined(__HIPCC__)
316 
317 inline Float8E4M3FNUZ operator"" _f8e4m3p8fnuz(unsigned long long int v) {
318  return Float8E4M3FNUZ(narrow<uint8_t>(v), Float8E4M3FNUZ::FromBits());
319 }
320 
321 inline Float8E4M3FNUZ operator"" _f8e4m3fnuzp8(long double v) {
322  return Float8E4M3FNUZ(static_cast<float>(v), true);
323 }
324 
325 #endif
326 
327 inline void Float8E4M3FNUZToFloat(const Float8E4M3FNUZ* blf, float* flt, size_t size) {
328  auto src = blf;
329  auto d = flt;
330  for (; size != 0; ++src, ++d, --size) {
331  *d = src->ToFloat();
332  }
333 }
334 
335 inline void FloatToFloat8E4M3FNUZ(const float* flt, Float8E4M3FNUZ* blf, size_t size, bool saturate) {
336  auto src = flt;
337  auto d = blf;
338  for (; size != 0; ++src, ++d, --size) {
339  new (d) Float8E4M3FNUZ(*src, saturate);
340  }
341 }
342 
343 // Float8E5M2
344 struct Float8E5M2 {
345  uint8_t val{0};
346 #if defined(__HIP__)
347  ORT_HOST_DEVICE Float8E5M2() = default;
348 #else
349  Float8E5M2() = default;
350 #endif
351 
352  struct FromBitsT {};
353  static constexpr ORT_HOST_DEVICE FromBitsT FromBits() { return FromBitsT(); }
354  constexpr ORT_HOST_DEVICE Float8E5M2(unsigned char bits, FromBitsT) : val(bits) {}
355 
356  inline explicit ORT_HOST_DEVICE Float8E5M2(float v, bool saturate = true) {
357 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
358  val = __nv_cvt_float_to_fp8(v, saturate ? __NV_SATFINITE : __NV_NOSAT, __NV_E5M2);
359 #else
360  uint32_t b;
361  std::memcpy(&b, &v, sizeof(b));
362 
363  val = (b & 0x80000000) >> 24; // sign
364  if ((b & 0x7FFFFFFF) == 0x7F800000) { // inf
365  if (saturate) {
366  // the highest available value
367  val |= 0x7B;
368  } else {
369  // the infinity
370  val |= 0x7C;
371  }
372  } else if ((b & 0x7F800000) == 0x7F800000) { // NaN
373  val |= 0x7f;
374  } else {
375  uint32_t e = (b & 0x7F800000) >> 23; // exponent
376  uint32_t m = b & 0x007FFFFF; // mantissa
377 
378  if (e != 0) {
379  if (e < 110) {
380  } else if (e < 113) {
381  // denormalized number
382  auto d = 112 - e;
383  if (d < 2) {
384  val |= 1 << (1 - d);
385  val |= m >> (22 + d);
386  } else if (m > 0) {
387  val |= 1;
388  }
389  auto mask = 1 << (21 + d);
390  if ((m & mask) && ((val & 1) || ((m & (mask - 1)) > 0) || ((m & mask) && (m & (mask << 1)) && ((m & (mask - 1)) == 0)))) {
391  // rounding
392  val += 1;
393  }
394  } else if (e < 143) { // 127 + 15 + 1
395  auto ex = e - 112; // 127 - 15
396  val |= ex << 2;
397  val |= m >> 21;
398  if ((m & 0x100000) && ((m & 0xFFFFF) || (m & 0x200000))) {
399  if ((val & 0x7F) < 0x7B) {
400  // rounding
401  val += 1;
402  } else if (saturate) {
403  val |= 0x7B;
404  } else {
405  val |= 0x7C;
406  }
407  }
408  } else if (saturate) {
409  val |= 0x7B;
410  } else {
411  val |= 0x7C;
412  }
413  }
414  }
415 #endif
416  }
417 
418  inline ORT_HOST_DEVICE float ToFloat() const {
419 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
420  return __half2float(__nv_cvt_fp8_to_halfraw(val, __NV_E5M2));
421 #else
422  uint32_t res;
423  if (val >= 253) {
424  res = 0xffc00000;
425  } else if (val >= 125 && val <= 127) {
426  res = 0x7fc00000;
427  } else if (val == 252) {
428  res = 0xff800000;
429  } else if (val == 124) {
430  res = 0x7f800000;
431  } else {
432  uint32_t expo = (val & 0x7C) >> 2;
433  uint32_t mant = val & 0x03;
434  uint32_t sign = val & 0x80;
435  res = sign << 24;
436  if (expo == 0) {
437  if (mant > 0) {
438  expo = 0x7F - 15;
439  if ((mant & 0x2) == 0) {
440  mant &= 0x1;
441  mant <<= 1;
442  expo -= 1;
443  }
444  res |= (mant & 0x1) << 22;
445  res |= expo << 23;
446  }
447  } else {
448  res |= mant << 21;
449  expo -= 15;
450  expo += 0x7F;
451  res |= expo << 23;
452  }
453  }
454 
455  float float_res;
456  std::memcpy(&float_res, &res, sizeof(float));
457  return float_res;
458 #endif
459  }
460 
461  inline ORT_HOST_DEVICE operator float() const { return ToFloat(); }
462 
463 #if defined(CUDA_VERSION) && CUDA_VERSION >= 11080
464  ORT_HOST_DEVICE Float8E5M2(const __nv_fp8_e5m2& value) { val = *reinterpret_cast<const unsigned char*>(&value); }
465  explicit ORT_HOST_DEVICE operator __nv_fp8_e5m2() const { return *reinterpret_cast<const __nv_fp8_e5m2*>(&val); }
466 #endif
467 };
468 
469 inline ORT_HOST_DEVICE bool operator==(const Float8E5M2& left, const Float8E5M2& right) { return left.val == right.val; }
470 inline ORT_HOST_DEVICE bool operator!=(const Float8E5M2& left, const Float8E5M2& right) { return left.val != right.val; }
471 inline ORT_HOST_DEVICE bool operator<(const Float8E5M2& left, const Float8E5M2& right) { return left.val < right.val; }
472 
473 // User defined suffixes to make it easier to declare
474 // initializers with MLFloat8E5M2 and Float8E5M2 from unsigned char
475 #if !defined(__CUDACC__) && !defined(__HIPCC__)
476 
477 inline Float8E5M2 operator"" _f8e5m2fn(unsigned long long int v) {
478  return Float8E5M2(narrow<uint8_t>(v), Float8E5M2::FromBits());
479 }
480 
481 inline Float8E5M2 operator"" _f8e5m2fnp8(long double v) {
482  return Float8E5M2(static_cast<float>(v), true);
483 }
484 
485 #endif
486 
487 inline void Float8E5M2ToFloat(const Float8E5M2* blf, float* flt, size_t size) {
488  auto src = blf;
489  auto d = flt;
490  for (; size != 0; ++src, ++d, --size) {
491  *d = src->ToFloat();
492  }
493 }
494 
495 inline void FloatToFloat8E5M2(const float* flt, Float8E5M2* blf, size_t size, bool saturate) {
496  auto src = flt;
497  auto d = blf;
498  for (; size != 0; ++src, ++d, --size) {
499  new (d) Float8E5M2(*src, saturate);
500  }
501 }
502 
503 // Float8E5M2FNUZ
505  uint8_t val{0};
506 #if defined(__HIP__)
507  ORT_HOST_DEVICE Float8E5M2FNUZ() = default;
508 #else
509  Float8E5M2FNUZ() = default;
510 #endif
511 
512  struct FromBitsT {};
513  static constexpr ORT_HOST_DEVICE FromBitsT FromBits() { return FromBitsT(); }
514  constexpr ORT_HOST_DEVICE Float8E5M2FNUZ(unsigned char bits, FromBitsT) : val(bits) {}
515 
516  inline explicit ORT_HOST_DEVICE Float8E5M2FNUZ(float v, bool saturate = true) {
517  // This type does not exist on CUDA.
518  uint32_t b;
519  std::memcpy(&b, &v, sizeof(b));
520 
521  val = (b & 0x80000000) >> 24; // sign
522  if ((b & 0x7FFFFFFF) == 0x7F800000) { // inf
523  if (saturate) {
524  val |= 0x7F;
525  } else {
526  val = 0x80;
527  }
528  } else if ((b & 0x7F800000) == 0x7F800000) { // NaN
529  val = 0x80;
530  } else {
531  uint32_t e = (b & 0x7F800000) >> 23; // exponent
532  uint32_t m = b & 0x007FFFFF; // mantissa
533 
534  if (e != 0) {
535  if (e < 109) {
536  } else if (e < 112) {
537  // denormalized number
538  auto d = 111 - e;
539  if (d < 2) {
540  val |= 1 << (1 - d);
541  val |= m >> (22 + d);
542  } else if (m > 0) {
543  val |= 1;
544  }
545  auto mask = 1 << (21 + d);
546  if ((m & mask) && ((val & 1) || ((m & (mask - 1)) > 0) || ((m & mask) && (m & (mask << 1)) && ((m & (mask - 1)) == 0)))) {
547  // rounding
548  val += 1;
549  }
550  } else if (e < 143) {
551  // normalized number
552  auto ex = e - 111;
553  val |= ex << 2;
554  val |= m >> 21;
555  if ((m & 0x100000) && ((m & 0xFFFFF) || (m & 0x200000))) {
556  if ((val & 0x7F) < 0x7F) {
557  // rounding
558  val += 1;
559  } else if (!saturate) {
560  val = 0x80;
561  }
562  }
563  } else if ((e == 255) && (m == 0)) {
564  val = 0x80;
565  } else if (saturate) {
566  val |= 0x7F;
567  } else {
568  val = 0x80;
569  }
570  } else if (m == 0) {
571  // -0
572  val = 0;
573  }
574  }
575  }
576 
577  inline ORT_HOST_DEVICE float ToFloat() const {
578  // This type does not exist on CUDA.
579  uint32_t res;
580  if (val == 0x80) {
581  res = 0xffc00000;
582  } else {
583  uint32_t expo = (val & 0x7C) >> 2;
584  uint32_t mant = val & 0x03;
585  uint32_t sign = val & 0x80;
586  res = sign << 24;
587  if (expo == 0) {
588  if (mant > 0) {
589  expo = 0x7F - 16;
590  if ((mant & 0x2) == 0) {
591  mant &= 0x1;
592  mant <<= 1;
593  expo -= 1;
594  }
595  res |= (mant & 0x1) << 22;
596  res |= expo << 23;
597  }
598  } else {
599  res |= mant << 21;
600  expo -= 16;
601  expo += 0x7F;
602  res |= expo << 23;
603  }
604  }
605 
606  float float_res;
607  std::memcpy(&float_res, &res, sizeof(float));
608  return float_res;
609  }
610 
611  inline ORT_HOST_DEVICE operator float() const { return ToFloat(); }
612 };
613 
614 inline ORT_HOST_DEVICE bool operator==(const Float8E5M2FNUZ& left, const Float8E5M2FNUZ& right) { return left.val == right.val; }
615 inline ORT_HOST_DEVICE bool operator!=(const Float8E5M2FNUZ& left, const Float8E5M2FNUZ& right) { return left.val != right.val; }
616 inline ORT_HOST_DEVICE bool operator<(const Float8E5M2FNUZ& left, const Float8E5M2FNUZ& right) { return left.val < right.val; }
617 
618 // User defined suffixes to make it easier to declare
619 // initializers with MLFloat8E5M2 and Float8E5M2 from unsigned char
620 #if !defined(__CUDACC__) && !defined(__HIPCC__)
621 
622 inline Float8E5M2FNUZ operator"" _f8e5m2fnuz(unsigned long long int v) {
623  return Float8E5M2FNUZ(narrow<uint8_t>(v), Float8E5M2FNUZ::FromBits());
624 }
625 
626 inline Float8E5M2FNUZ operator"" _f8e5m2fnuzp8(long double v) {
627  return Float8E5M2FNUZ(static_cast<float>(v), true);
628 }
629 
630 #endif
631 
632 inline void Float8E5M2FNUZToFloat(const Float8E5M2FNUZ* blf, float* flt, size_t size) {
633  auto src = blf;
634  auto d = flt;
635  for (; size != 0; ++src, ++d, --size) {
636  *d = src->ToFloat();
637  }
638 }
639 
640 inline void FloatToFloat8E5M2FNUZ(const float* flt, Float8E5M2FNUZ* blf, size_t size, bool saturate) {
641  auto src = flt;
642  auto d = blf;
643  for (; size != 0; ++src, ++d, --size) {
644  new (d) Float8E5M2FNUZ(*src, saturate);
645  }
646 }
647 
648 } // namespace onnxruntime
649 
650 #endif // DISABLE_FLOAT8_TYPES
void FloatToFloat8E4M3FNUZ(const float *flt, Float8E4M3FNUZ *blf, size_t size, bool saturate)
Definition: float8.h:335
static constexpr ORT_HOST_DEVICE FromBitsT FromBits()
Definition: float8.h:36
GLint left
Definition: glcorearb.h:2005
const GLdouble * v
Definition: glcorearb.h:837
GLsizei const GLfloat * value
Definition: glcorearb.h:824
ORT_HOST_DEVICE bool operator<(const Float8E4M3FN &left, const Float8E4M3FN &right)
Definition: float8.h:158
GLdouble right
Definition: glad.h:2817
ORT_HOST_DEVICE float ToFloat() const
Definition: float8.h:577
ORT_HOST_DEVICE float ToFloat() const
Definition: float8.h:418
void Float8E4M3FNUZToFloat(const Float8E4M3FNUZ *blf, float *flt, size_t size)
Definition: float8.h:327
ORT_HOST_DEVICE Float8E4M3FNUZ(float v, bool saturate=true)
Definition: float8.h:203
static constexpr ORT_HOST_DEVICE FromBitsT FromBits()
Definition: float8.h:513
static constexpr ORT_HOST_DEVICE FromBitsT FromBits()
Definition: float8.h:353
constexpr ORT_HOST_DEVICE Float8E4M3FN(unsigned char bits, FromBitsT)
Definition: float8.h:37
ORT_HOST_DEVICE bool operator==(const Float8E4M3FN &left, const Float8E4M3FN &right)
Definition: float8.h:156
GLdouble GLdouble x2
Definition: glad.h:2349
ORT_HOST_DEVICE bool operator!=(const Float8E4M3FN &left, const Float8E4M3FN &right)
Definition: float8.h:157
void FloatToFloat8E4M3FN(const float *flt, Float8E4M3FN *blf, size_t size, bool saturate)
Definition: float8.h:182
ORT_HOST_DEVICE Float8E5M2FNUZ(float v, bool saturate=true)
Definition: float8.h:516
IMATH_NAMESPACE::V2f float
GLint GLuint mask
Definition: glcorearb.h:124
constexpr ORT_HOST_DEVICE Float8E4M3FNUZ(unsigned char bits, FromBitsT)
Definition: float8.h:201
ORT_HOST_DEVICE Float8E5M2(float v, bool saturate=true)
Definition: float8.h:356
ImageBuf OIIO_API saturate(const ImageBuf &src, float scale=0.0f, int firstchannel=0, ROI roi={}, int nthreads=0)
IMATH_HOSTDEVICE constexpr int sign(T a) IMATH_NOEXCEPT
Definition: ImathFun.h:33
#define ORT_HOST_DEVICE
Definition: float8.h:24
GLboolean GLboolean GLboolean b
Definition: glcorearb.h:1222
static constexpr ORT_HOST_DEVICE FromBitsT FromBits()
Definition: float8.h:200
void FloatToFloat8E5M2(const float *flt, Float8E5M2 *blf, size_t size, bool saturate)
Definition: float8.h:495
ORT_HOST_DEVICE Float8E4M3FN(float v, bool saturate=true)
Definition: float8.h:39
GLsizeiptr size
Definition: glcorearb.h:664
void Float8E5M2ToFloat(const Float8E5M2 *blf, float *flt, size_t size)
Definition: float8.h:487
void FloatToFloat8E5M2FNUZ(const float *flt, Float8E5M2FNUZ *blf, size_t size, bool saturate)
Definition: float8.h:640
ORT_HOST_DEVICE float ToFloat() const
Definition: float8.h:268
GLuint GLfloat * val
Definition: glcorearb.h:1608
constexpr ORT_HOST_DEVICE Float8E5M2(unsigned char bits, FromBitsT)
Definition: float8.h:354
void Float8E4M3FNToFloat(const Float8E4M3FN *blf, float *flt, size_t size)
Definition: float8.h:174
Definition: core.h:1131
void Float8E5M2FNUZToFloat(const Float8E5M2FNUZ *blf, float *flt, size_t size)
Definition: float8.h:632
GLenum src
Definition: glcorearb.h:1793
constexpr ORT_HOST_DEVICE Float8E5M2FNUZ(unsigned char bits, FromBitsT)
Definition: float8.h:514
ORT_HOST_DEVICE float ToFloat() const
Definition: float8.h:105