diff --git a/src/api/include/migraphx/migraphx.h b/src/api/include/migraphx/migraphx.h index 90ba7c3e017..cdc5bf3e9a5 100644 --- a/src/api/include/migraphx/migraphx.h +++ b/src/api/include/migraphx/migraphx.h @@ -47,7 +47,8 @@ m(uint64_type, uint64_t) \ m(fp8e4m3fnuz_type, migraphx::fp8::fp8e4m3fnuz) \ m(fp8e4m3fn_type, migraphx::fp8::fp8e4m3fn) \ - m(fp8e5m2_type, migraphx::fp8::fp8e5m2) + m(fp8e5m2_type, migraphx::fp8::fp8e5m2) \ + m(bf16_type, bf16) // clang-format on #ifdef __cplusplus diff --git a/src/driver/main.cpp b/src/driver/main.cpp index 04fa0cfe3bc..f193f1a9a07 100644 --- a/src/driver/main.cpp +++ b/src/driver/main.cpp @@ -482,6 +482,7 @@ struct compiler compiler_target ct; compile_options co; bool to_fp16 = false; + bool to_bf16 = false; bool to_fp8 = false; bool to_int8 = false; bool to_int4 = false; @@ -506,9 +507,11 @@ struct compiler ap.help("Exhastively search for best tuning parameters for kernels"), ap.set_value(true)); ap(to_fp16, {"--fp16"}, ap.help("Quantize for fp16"), ap.set_value(true)); + ap(to_bf16, {"--bf16"}, ap.help("Quantize for bf16"), ap.set_value(true)); ap(to_int8, {"--int8"}, ap.help("Quantize for int8"), ap.set_value(true)); ap(to_fp8, {"--fp8"}, ap.help("Quantize for fp8"), ap.set_value(true)); ap(to_int4, {"--int4-weights"}, ap.help("Quantize weights for int4"), ap.set_value(true)); + ap(to_bf16, {"--bf16"}, ap.help("Quantize for fp16"), ap.set_value(true)); } auto params(const program& p) @@ -555,6 +558,10 @@ struct compiler { quantize_fp16(p); } + if(to_bf16) + { + quantize_bf16(p); + } if(to_int8) { quantize_int8(p, t, {host_params(p)}); @@ -567,6 +574,10 @@ struct compiler { quantize_int4_weights(p); } + if(to_bf16) + { + quantize_bf16(p); + } p.compile(t, co); l.save(p); return p; @@ -639,6 +650,10 @@ struct verify : command { vo.quantize = precision::fp16; } + if(c.to_bf16) + { + vo.quantize = precision::bf16; + } if(c.to_int8) { vo.quantize = precision::int8; diff --git a/src/driver/precision.hpp b/src/driver/precision.hpp index d7d7cecf00e..9ed1f402f9d 100644 --- a/src/driver/precision.hpp +++ b/src/driver/precision.hpp @@ -32,6 +32,7 @@ enum class precision { fp32, fp16, + bf16, int8 }; diff --git a/src/driver/verify.cpp b/src/driver/verify.cpp index 92bae3eee86..599c5034565 100644 --- a/src/driver/verify.cpp +++ b/src/driver/verify.cpp @@ -100,6 +100,10 @@ std::vector run_target(program p, { quantize_fp16(p); } + if(vo.quantize == precision::bf16) + { + quantize_bf16(p); + } p.compile(t, options); parameter_map m; diff --git a/src/include/migraphx/bf16.hpp b/src/include/migraphx/bf16.hpp new file mode 100644 index 00000000000..26ecdd7c996 --- /dev/null +++ b/src/include/migraphx/bf16.hpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#ifndef MIGRAPHX_GUARD_RTGLIB_BF16_HPP +#define MIGRAPHX_GUARD_RTGLIB_BF16_HPP + +#include +#include + +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { + +using bf16 = migraphx::generic_float<7, 8>; + +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx + +#endif diff --git a/src/include/migraphx/generic_float.hpp b/src/include/migraphx/generic_float.hpp index a09c19a3a26..7ff18615904 100644 --- a/src/include/migraphx/generic_float.hpp +++ b/src/include/migraphx/generic_float.hpp @@ -104,6 +104,8 @@ struct float32_parts unsigned int exponent : 8; unsigned int sign : 1; + static constexpr unsigned int exponent_width() { return 8; } + static constexpr unsigned int mantissa_width() { return 23; } static constexpr unsigned int max_exponent() { return all_ones<8>(); } @@ -152,7 +154,7 @@ struct __attribute__((packed, may_alias)) generic_float float32_parts f{}; f.sign = sign; - if(exponent == 0) // subnormal fps + if(exponent == 0 and ExponentSize != float32_parts::exponent_width()) // subnormal fps { if(mantissa == 0) diff --git a/src/include/migraphx/quantization.hpp b/src/include/migraphx/quantization.hpp index d849023b6cf..eead5e40ba1 100644 --- a/src/include/migraphx/quantization.hpp +++ b/src/include/migraphx/quantization.hpp @@ -51,6 +51,9 @@ quantize_fp8(program& prog, const target& t, const std::vector& c MIGRAPHX_EXPORT void quantize_int4_weights(program& prog); +MIGRAPHX_EXPORT void quantize_bf16(program& prog, + const std::vector& ins_names = {"all"}); + } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/include/migraphx/shape.hpp b/src/include/migraphx/shape.hpp index e15c4dece44..d5004f63bc9 100644 --- a/src/include/migraphx/shape.hpp +++ b/src/include/migraphx/shape.hpp @@ -34,6 +34,7 @@ #include #include #include +#include #include #include #include @@ -64,8 +65,9 @@ struct MIGRAPHX_EXPORT shape m(uint64_type, uint64_t) \ m(fp8e4m3fnuz_type, migraphx::fp8::fp8e4m3fnuz) \ m(fp8e4m3fn_type, migraphx::fp8::fp8e4m3fn) \ - m(fp8e5m2_type, migraphx::fp8::fp8e5m2) -// clang-format on + m(fp8e5m2_type, migraphx::fp8::fp8e5m2) \ + m(bf16_type, bf16) + // clang-format on #define MIGRAPHX_SHAPE_GENERATE_ENUM_TYPES(x, t) x, enum type_t diff --git a/src/include/migraphx/type_traits.hpp b/src/include/migraphx/type_traits.hpp index 908f67d9f14..ff6f9225d6a 100644 --- a/src/include/migraphx/type_traits.hpp +++ b/src/include/migraphx/type_traits.hpp @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -53,6 +54,10 @@ MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_floating_point, half) MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_signed, half) MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_arithmetic, half) +MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_floating_point, bf16) +MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_signed, bf16) +MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_arithmetic, bf16) + MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_floating_point, migraphx::fp8::fp8e4m3fnuz) MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_signed, migraphx::fp8::fp8e4m3fnuz) MIGRAPHX_DETAIL_EXTEND_TRAIT_FOR(is_arithmetic, migraphx::fp8::fp8e4m3fnuz) diff --git a/src/py/migraphx_py.cpp b/src/py/migraphx_py.cpp index 9d05e32e67d..479a0289919 100644 --- a/src/py/migraphx_py.cpp +++ b/src/py/migraphx_py.cpp @@ -180,6 +180,17 @@ struct npy_format_descriptor static constexpr auto name() { return _("fp8e5m2"); } }; +template <> +struct npy_format_descriptor +{ + static std::string format() + { + // TODO: no standard format in numpy for bf16 + return "z"; + } + static constexpr auto name() { return _("bf16"); } +}; + } // namespace detail } // namespace pybind11 @@ -623,6 +634,10 @@ MIGRAPHX_PYBIND11_MODULE(migraphx, m) }, "Auto-convert FP8 parameters and return values to Float for MIGraphX Program", py::arg("prog")); + m.def("quantize_bf16", + &migraphx::quantize_bf16, + py::arg("prog"), + py::arg("ins_names") = std::vector{"all"}); #ifdef HAVE_GPU m.def("allocate_gpu", &migraphx::gpu::allocate_gpu, py::arg("s"), py::arg("host") = false); diff --git a/src/quantization.cpp b/src/quantization.cpp index 7e02ae66685..276012bbf73 100644 --- a/src/quantization.cpp +++ b/src/quantization.cpp @@ -74,6 +74,16 @@ void quantize_fp16(program& prog, const std::vector& ins_names) quant_tracer()); } +void quantize_bf16(program& prog, const std::vector& ins_names) +{ + run_passes(prog, + {normalize_ops{}, + optimize_module{{"quantizelinear", "dequantizelinear"}}, + truncate_float_pass{ins_names, shape::bf16_type}, + optimize_module{{"quantizelinear", "dequantizelinear"}}}, + quant_tracer()); +} + void quantize_8bits(program& prog, const target& t, shape::type_t precision, diff --git a/src/targets/gpu/fuse_mlir.cpp b/src/targets/gpu/fuse_mlir.cpp index 4376bb323cc..5848ff64f77 100644 --- a/src/targets/gpu/fuse_mlir.cpp +++ b/src/targets/gpu/fuse_mlir.cpp @@ -361,6 +361,7 @@ bool is_pointwise_op_supported_by_mlir(const instruction& i) const auto& name = i.name(); const auto result_type = i.get_shape().type(); const std::initializer_list allowed_types = {type_t::float_type, + type_t::bf16_type, type_t::half_type, type_t::fp8e4m3fnuz_type, type_t::fp8e4m3fn_type, @@ -407,6 +408,7 @@ bool is_pointwise_op_supported_by_mlir(const instruction& i) }; std::set float_types = {type_t::float_type, type_t::half_type, + type_t::bf16_type, type_t::fp8e4m3fnuz_type, type_t::fp8e4m3fn_type, type_t::fp8e5m2_type}; @@ -426,7 +428,7 @@ bool is_pointwise_op_supported_by_mlir(const instruction& i) return false; } // else return std::all_of(i.inputs().begin(), i.inputs().end(), [](const auto& arg) { - return contains({type_t::float_type, type_t::half_type}, arg->get_shape().type()); + return contains({type_t::float_type, type_t::half_type, type_t::bf16_type}, arg->get_shape().type()); }); } return false; @@ -438,7 +440,7 @@ bool is_reduce_op_supported_by_mlir(const instruction& i) const auto& name = i.name(); const auto result_type = i.get_shape().type(); const std::initializer_list allowed_types = { - type_t::float_type, type_t::half_type, type_t::fp8e4m3fnuz_type}; + type_t::float_type, type_t::half_type, type_t::bf16_type, type_t::fp8e4m3fnuz_type}; // Preliminary type check. if(not contains(allowed_types, result_type)) { @@ -695,6 +697,7 @@ struct find_mlir_standalone_op if(std::any_of(gemm_based_op->inputs().begin(), gemm_based_op->inputs().end(), [&](auto i) { return not contains({shape::type_t::float_type, shape::type_t::half_type, + shape::type_t::bf16_type, shape::type_t::int8_type, shape::type_t::fp8e4m3fnuz_type, shape::type_t::fp8e4m3fn_type, diff --git a/src/targets/gpu/gemm_impl.cpp b/src/targets/gpu/gemm_impl.cpp index 19d4f056deb..caaaf53d35e 100644 --- a/src/targets/gpu/gemm_impl.cpp +++ b/src/targets/gpu/gemm_impl.cpp @@ -72,6 +72,7 @@ rocblas_datatype get_type(shape::type_t type) case shape::int16_type: case shape::int64_type: case shape::uint64_type: MIGRAPHX_THROW("ROCBLAS_GEMM: data type not supported!"); + case shape::bf16_type: return rocblas_datatype_bf16_r; } MIGRAPHX_THROW("ROCBLAS_GEMM: data type not supported!"); @@ -221,7 +222,7 @@ struct gemm_impl compute_type = rb_compute_type{output_type}; if(compute_fp32) { - if(arg_type == rocblas_datatype_f16_r) + if(arg_type == rocblas_datatype_f16_r or arg_type == rocblas_datatype_bf16_r) compute_type = rocblas_datatype_f32_r; } if(arg_type == rocblas_datatype_f8_r) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index e4a3c256563..9b95e466107 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -84,6 +84,7 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::int16_type: case shape::int64_type: case shape::uint64_type: MIGRAPHX_THROW("HIPBLAS_GEMM: data type not supported!"); + case shape::bf16_type: return HIP_R_16BF; } MIGRAPHX_THROW("HIPBLAS_GEMM: data type not supported!"); diff --git a/src/targets/gpu/kernels/include/migraphx/kernels/hip.hpp b/src/targets/gpu/kernels/include/migraphx/kernels/hip.hpp index 8ddc7ad0ef6..5e1c645d784 100644 --- a/src/targets/gpu/kernels/include/migraphx/kernels/hip.hpp +++ b/src/targets/gpu/kernels/include/migraphx/kernels/hip.hpp @@ -27,7 +27,9 @@ #ifndef MIGRAPHX_USE_HIPRTC #include #include +#include #include +#include #endif #endif // MIGRAPHX_GUARD_KERNELS_HIP_HPP diff --git a/src/targets/gpu/kernels/include/migraphx/kernels/types.hpp b/src/targets/gpu/kernels/include/migraphx/kernels/types.hpp index f65cdfbba34..55a0cc5c2a6 100644 --- a/src/targets/gpu/kernels/include/migraphx/kernels/types.hpp +++ b/src/targets/gpu/kernels/include/migraphx/kernels/types.hpp @@ -76,6 +76,7 @@ using vec = T __attribute__((ext_vector_type(N))); using half = _Float16; using half2 = migraphx::vec; +using bf16 = __bf16; } // namespace migraphx diff --git a/src/targets/gpu/mlir.cpp b/src/targets/gpu/mlir.cpp index e96c6a5d7dd..75d968055d5 100644 --- a/src/targets/gpu/mlir.cpp +++ b/src/targets/gpu/mlir.cpp @@ -312,6 +312,8 @@ struct mlir_program result = mlirF32TypeGet(ctx.get()); else if(as.type_enum() == shape::half_type) result = mlirF16TypeGet(ctx.get()); + else if(as.type_enum() == shape::bf16_type) + result = mlirBF16TypeGet(ctx.get()); else if(as.type_enum() == shape::fp8e4m3fnuz_type) result = mlirFloat8E4M3FNUZTypeGet(ctx.get()); else if(as.type_enum() == shape::fp8e4m3fn_type) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index fa01c514c0c..f0e7740e4fa 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -99,6 +99,7 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti unsupported_types.erase(shape::type_t::uint8_type); unsupported_types.erase(shape::type_t::int32_type); unsupported_types.erase(shape::type_t::tuple_type); + unsupported_types.erase(shape::type_t::bf16_type); // whiltelist supported Ops for the FP8 types // different between fp8e4m3fnuz and OCP types because rocBLAS only has diff --git a/test/bf16.cpp b/test/bf16.cpp new file mode 100644 index 00000000000..9ba04d083a2 --- /dev/null +++ b/test/bf16.cpp @@ -0,0 +1,1245 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ +#include +#include +#include +#include +#include "test.hpp" + +#include +#include +#include +#include +#include +#include + +template +bool bit_equal(const T& x, const U& y) +{ + static_assert(sizeof(T) == sizeof(U)); + using type = std::array; + return migraphx::bit_cast(x) == migraphx::bit_cast(y); +} + +TEST_CASE(check_numeric_limits) +{ + CHECK(bit_equal(std::numeric_limits::min(), uint16_t{0x0080})); + CHECK(bit_equal(std::numeric_limits::lowest(), uint16_t{0xff7f})); + CHECK(bit_equal(std::numeric_limits::max(), uint16_t{0x7f7f})); + CHECK(bit_equal(std::numeric_limits::epsilon(), uint16_t{0x3c00})); + CHECK(bit_equal(std::numeric_limits::denorm_min(), uint16_t{0x0001})); + CHECK(bit_equal(std::numeric_limits::infinity(), uint16_t{0x7f80})); + CHECK(bit_equal(std::numeric_limits::quiet_NaN(), uint16_t{0x7fc0})); + CHECK(bit_equal(std::numeric_limits::signaling_NaN(), uint16_t{0x7fa0})); +} + +const std::map& bf16_lut() // NOLINT(readability-function-size) +{ + static const std::map result = { + {0x0000, 0.0}, + {0x002d, 0.00000000000000000000000000000000000000413259732711}, + {0x004e, 0.00000000000000000000000000000000000000716316870032}, + {0x005b, 0.00000000000000000000000000000000000000835703015038}, + {0x00cd, 0.00000000000000000000000000000000000001882627671239}, + {0x00ce, 0.00000000000000000000000000000000000001891811220855}, + {0x0170, 0.00000000000000000000000000000000000004408103815584}, + {0x01be, 0.00000000000000000000000000000000000006979497708007}, + {0x01fe, 0.00000000000000000000000000000000000009330486409652}, + {0x0211, 0.00000000000000000000000000000000000010652917554327}, + {0x028f, 0.00000000000000000000000000000000000021011961520948}, + {0x02bc, 0.00000000000000000000000000000000000027624117244324}, + {0x02e8, 0.00000000000000000000000000000000000034089336173846}, + {0x039c, 0.00000000000000000000000000000000000091688559364138}, + {0x03a9, 0.00000000000000000000000000000000000099329272644483}, + {0x03cf, 0.00000000000000000000000000000000000121663665310107}, + {0x03da, 0.00000000000000000000000000000000000128128884239629}, + {0x03f3, 0.00000000000000000000000000000000000142822563624908}, + {0x0427, 0.00000000000000000000000000000000000196307556587322}, + {0x044c, 0.00000000000000000000000000000000000239800847567747}, + {0x0485, 0.00000000000000000000000000000000000312681497318728}, + {0x0498, 0.00000000000000000000000000000000000357350282649975}, + {0x04da, 0.00000000000000000000000000000000000512515536958517}, + {0x051b, 0.00000000000000000000000000000000000728806497509818}, + {0x0533, 0.00000000000000000000000000000000000841653955188758}, + {0x0536, 0.00000000000000000000000000000000000855759887398625}, + {0x0577, 0.00000000000000000000000000000000001161388418612420}, + {0x0587, 0.00000000000000000000000000000000001269533898888071}, + {0x065a, 0.00000000000000000000000000000000004100124295668139}, + {0x0735, 0.00000000000000000000000000000000013616926559925378}, + {0x075e, 0.00000000000000000000000000000000016701423736483061}, + {0x0765, 0.00000000000000000000000000000000017228045205651446}, + {0x076b, 0.00000000000000000000000000000000017679435036367204}, + {0x07ba, 0.00000000000000000000000000000000027986169504377021}, + {0x07d1, 0.00000000000000000000000000000000031446824873197835}, + {0x07db, 0.00000000000000000000000000000000032951457642250363}, + {0x07eb, 0.00000000000000000000000000000000035358870072734408}, + {0x080d, 0.00000000000000000000000000000000042430644087281290}, + {0x0841, 0.00000000000000000000000000000000058078824885427581}, + {0x08ff, 0.00000000000000000000000000000000153472542443357857}, + {0x09e1, 0.00000000000000000000000000000000541667796858910084}, + {0x0a0b, 0.00000000000000000000000000000000669260655674564459}, + {0x0a0f, 0.00000000000000000000000000000000688519955118436817}, + {0x0a12, 0.00000000000000000000000000000000702964429701341086}, + {0x0a1d, 0.00000000000000000000000000000000755927503171990072}, + {0x0aa2, 0.00000000000000000000000000000001560003254953661041}, + {0x0aaf, 0.00000000000000000000000000000001685188701338831371}, + {0x0ab4, 0.00000000000000000000000000000001733336949948512268}, + {0x0ade, 0.00000000000000000000000000000002137782238269831797}, + {0x0b6e, 0.00000000000000000000000000000004583713267641621330}, + {0x0bb6, 0.00000000000000000000000000000007010384997569538505}, + {0x0c8c, 0.00000000000000000000000000000021570415377137041554}, + {0x0cf7, 0.00000000000000000000000000000038056375701091780456}, + {0x0d06, 0.00000000000000000000000000000041291938007662336690}, + {0x0d20, 0.00000000000000000000000000000049303806576313237838}, + {0x0d5e, 0.00000000000000000000000000000068409031624634617501}, + {0x0d60, 0.00000000000000000000000000000069025329206838532974}, + {0x0d6a, 0.00000000000000000000000000000072106817117858110338}, + {0x0d76, 0.00000000000000000000000000000075804602611081603176}, + {0x0d7e, 0.00000000000000000000000000000078269792939897265068}, + {0x0dcc, 0.00000000000000000000000000000125724706769598756487}, + {0x0dfc, 0.00000000000000000000000000000155306990715386699190}, + {0x0e09, 0.00000000000000000000000000000168865537523872839596}, + {0x0e3c, 0.00000000000000000000000000000231727890908672217840}, + {0x0e69, 0.00000000000000000000000000000287194673307024610408}, + {0x0f42, 0.00000000000000000000000000000956493847580476814062}, + {0x0f8f, 0.00000000000000000000000000001410088868082558602173}, + {0x0fa2, 0.00000000000000000000000000001597443333072548905959}, + {0x1007, 0.00000000000000000000000000002662405555120914843265}, + {0x1015, 0.00000000000000000000000000002938506871948268975159}, + {0x1030, 0.00000000000000000000000000003470987982972451943812}, + {0x1042, 0.00000000000000000000000000003825975390321907256247}, + {0x1056, 0.00000000000000000000000000004220405842932413158953}, + {0x1166, 0.00000000000000000000000000018143800820083271524470}, + {0x1182, 0.00000000000000000000000000020510383535746306940705}, + {0x128a, 0.00000000000000000000000000087090243936399703317455}, + {0x1295, 0.00000000000000000000000000094032219902344607205078}, + {0x129e, 0.00000000000000000000000000099712018419935892204042}, + {0x12b0, 0.00000000000000000000000000111071615455118462201971}, + {0x12bf, 0.00000000000000000000000000120537946317770603866912}, + {0x1315, 0.00000000000000000000000000188064439804689214410156}, + {0x1343, 0.00000000000000000000000000246124602428955683288459}, + {0x13e2, 0.00000000000000000000000000570504206655835737673762}, + {0x13ee, 0.00000000000000000000000000600796465416322591001572}, + {0x1445, 0.00000000000000000000000000994595829302651684263107}, + {0x1469, 0.00000000000000000000000001176349381865572804229970}, + {0x14bd, 0.00000000000000000000000001908412301910671759652054}, + {0x1541, 0.00000000000000000000000003897603960515975128178268}, + {0x1546, 0.00000000000000000000000003998578156384264639270970}, + {0x1569, 0.00000000000000000000000004705397527462291216919879}, + {0x16a1, 0.00000000000000000000000026010952855671378057479844}, + {0x16b6, 0.00000000000000000000000029403685836845905630194606}, + {0x16f4, 0.00000000000000000000000039420326066980225130590570}, + {0x1703, 0.00000000000000000000000042328382907986963050060367}, + {0x1720, 0.00000000000000000000000051698788284564229679463043}, + {0x178e, 0.00000000000000000000000091765349205101507681046902}, + {0x17f6, 0.00000000000000000000000158973773975035006264348858}, + {0x18e8, 0.00000000000000000000000599705944100945064281771302}, + {0x18f1, 0.00000000000000000000000622970398828998967637529671}, + {0x1927, 0.00000000000000000000000863369764352222635647032822}, + {0x1a39, 0.00000000000000000000003825710333057752996280265201}, + {0x1a60, 0.00000000000000000000004632211430296954979279888676}, + {0x1a69, 0.00000000000000000000004818327068121386206125955631}, + {0x1b1f, 0.00000000000000000000013152171739593140030455398204}, + {0x1b43, 0.00000000000000000000016130021944784039659992469495}, + {0x1ba0, 0.00000000000000000000026469779601696885595885078146}, + {0x1bc0, 0.00000000000000000000031763735522036262715062093775}, + {0x1bea, 0.00000000000000000000038712052667481695183981926789}, + {0x1c25, 0.00000000000000000000054593920428499826541512973677}, + {0x1c32, 0.00000000000000000000058895259613775570450844298875}, + {0x1cbe, 0.00000000000000000000125731453108060206580454121195}, + {0x1ccd, 0.00000000000000000000135657620458696538678911025499}, + {0x1cd6, 0.00000000000000000000141613320869078337937985168082}, + {0x1d23, 0.00000000000000000000215728703753829617606463386892}, + {0x1d38, 0.00000000000000000000243521972335611347482142718945}, + {0x1dce, 0.00000000000000000000545277459794955843275232609813}, + {0x1e7c, 0.00000000000000000001334076891925523034032607938570}, + {0x1e9e, 0.00000000000000000001672890070827243169659936938842}, + {0x1eb1, 0.00000000000000000001874060395800139500188663532754}, + {0x1f33, 0.00000000000000000003790472438962994017330743190541}, + {0x1f56, 0.00000000000000000004531626267810506814015525378636}, + {0x1fa8, 0.00000000000000000007115076756936122848173909005709}, + {0x1fed, 0.00000000000000000010037340424963459017959621633054}, + {0x2001, 0.00000000000000000010926725019580474373981360258767}, + {0x213f, 0.00000000000000000064713317170228545904819839051925}, + {0x2154, 0.00000000000000000071828393927164668752993748057634}, + {0x216b, 0.00000000000000000079621097041904231872422315063886}, + {0x219f, 0.00000000000000000107742590890747003129490622086450}, + {0x2225, 0.00000000000000000223616698075135289514037140179425}, + {0x2229, 0.00000000000000000229037708937562811684074404183775}, + {0x2274, 0.00000000000000000330681662608078852372273104265332}, + {0x227b, 0.00000000000000000340168431617327016169838316272944}, + {0x2294, 0.00000000000000000401154803819636640582757536321878}, + {0x22b7, 0.00000000000000000496022493912118278558409656397998}, + {0x2379, 0.00000000000000001349831704744453020339278737083077}, + {0x2389, 0.00000000000000001485356976305141074590210337191820}, + {0x2464, 0.00000000000000004943961906533900219073984771966934}, + {0x2475, 0.00000000000000005312590645178971726636518724262714}, + {0x2491, 0.00000000000000006288372600415925717243226245045662}, + {0x24f6, 0.00000000000000010668549377257363630633335560560226}, + {0x2567, 0.00000000000000020036056147532121940457727760076523}, + {0x256a, 0.00000000000000020296264668928643004619516432285309}, + {0x257f, 0.00000000000000022117724318704290453752037137746811}, + {0x25f9, 0.00000000000000043194614551822496650856919586658478}, + {0x263b, 0.00000000000000064878658001532585331005975604057312}, + {0x270c, 0.00000000000000194289029309402394574135541915893555}, + {0x272c, 0.00000000000000238697950294408656191080808639526367}, + {0x2745, 0.00000000000000273392419813944798079319298267364502}, + {0x2791, 0.00000000000000402455846426619245903566479682922363}, + {0x2792, 0.00000000000000405231403988182137254625558853149414}, + {0x27a0, 0.00000000000000444089209850062616169452667236328125}, + {0x27a3, 0.00000000000000452415882534751290222629904747009277}, + {0x27d4, 0.00000000000000588418203051332966424524784088134766}, + {0x27dd, 0.00000000000000613398221105398988584056496620178223}, + {0x2821, 0.00000000000000893729534823251015041023492813110352}, + {0x28d9, 0.00000000000002409183963436589692719280719757080078}, + {0x2981, 0.00000000000005728750807065807748585939407348632812}, + {0x29ca, 0.00000000000008970602038971264846622943878173828125}, + {0x2a5b, 0.00000000000019451107391432742588222026824951171875}, + {0x2aa8, 0.00000000000029842794901924207806587219238281250}, + {0x2ac4, 0.000000000000348165940522449091076850891113281250}, + {0x2ae7, 0.00000000000041033842990145785734057426452636718750}, + {0x2af1, 0.00000000000042810199829546036198735237121582031250}, + {0x2afe, 0.0000000000004511946372076636180281639099121093750}, + {0x2b24, 0.00000000000058264504332328215241432189941406250}, + {0x2b4c, 0.00000000000072475359047530218958854675292968750}, + {0x2b85, 0.000000000000945021838560933247208595275878906250}, + {0x2bed, 0.000000000001683986283751437440514564514160156250}, + {0x2c18, 0.00000000000216004991671070456504821777343750}, + {0x2cdf, 0.0000000000063380412029800936579704284667968750}, + {0x2d6b, 0.000000000013358203432289883494377136230468750}, + {0x2d96, 0.0000000000170530256582424044609069824218750}, + {0x2da2, 0.0000000000184172677109017968177795410156250}, + {0x2db8, 0.00000000002091837814077734947204589843750}, + {0x2de1, 0.00000000002557953848736360669136047363281250}, + {0x2e90, 0.00000000006548361852765083312988281250}, + {0x2ea3, 0.000000000074123818194493651390075683593750}, + {0x2ef0, 0.00000000010913936421275138854980468750}, + {0x2f09, 0.00000000012460077414289116859436035156250}, + {0x2f6b, 0.00000000021373125491663813591003417968750}, + {0x303b, 0.000000000680302036926150321960449218750}, + {0x308f, 0.00000000104046193882822990417480468750}, + {0x309c, 0.000000001135049387812614440917968750}, + {0x30a9, 0.00000000122963683679699897766113281250}, + {0x312a, 0.000000002473825588822364807128906250}, + {0x313d, 0.0000000027503119781613349914550781250}, + {0x3159, 0.0000000031577656045556068420410156250}, + {0x31c6, 0.00000000576255843043327331542968750}, + {0x3212, 0.0000000084983184933662414550781250}, + {0x3245, 0.00000001146690919995307922363281250}, + {0x329b, 0.0000000180443748831748962402343750}, + {0x32ba, 0.000000021653249859809875488281250}, + {0x32cc, 0.00000002374872565269470214843750}, + {0x3332, 0.00000004144385457038879394531250}, + {0x33c4, 0.000000091269612312316894531250}, + {0x3424, 0.00000015273690223693847656250}, + {0x3589, 0.0000010207295417785644531250}, + {0x3594, 0.00000110268592834472656250}, + {0x368b, 0.00000414252281188964843750}, + {0x36a0, 0.000004768371582031250}, + {0x36e9, 0.00000694394111633300781250}, + {0x36ed, 0.00000706315040588378906250}, + {0x3750, 0.000012397766113281250}, + {0x375f, 0.0000132918357849121093750}, + {0x37ce, 0.00002455711364746093750}, + {0x37d2, 0.00002503395080566406250}, + {0x37f0, 0.00002861022949218750}, + {0x380e, 0.0000338554382324218750}, + {0x3826, 0.0000395774841308593750}, + {0x387f, 0.00006079673767089843750}, + {0x38e1, 0.0001072883605957031250}, + {0x38e9, 0.0001111030578613281250}, + {0x3964, 0.0002174377441406250}, + {0x3994, 0.000282287597656250}, + {0x3a26, 0.000633239746093750}, + {0x3a2c, 0.00065612792968750}, + {0x3a6a, 0.000892639160156250}, + {0x3a85, 0.001014709472656250}, + {0x3ab9, 0.001411437988281250}, + {0x3aba, 0.00141906738281250}, + {0x3af7, 0.001884460449218750}, + {0x3b03, 0.00199890136718750}, + {0x3bb3, 0.0054626464843750}, + {0x3bbf, 0.0058288574218750}, + {0x3be8, 0.0070800781250}, + {0x3c06, 0.00817871093750}, + {0x3c29, 0.010314941406250}, + {0x3c3f, 0.011657714843750}, + {0x3c73, 0.014831542968750}, + {0x3ce9, 0.02844238281250}, + {0x3cfa, 0.0305175781250}, + {0x3cfb, 0.03063964843750}, + {0x3d0f, 0.0349121093750}, + {0x3d2a, 0.041503906250}, + {0x3d43, 0.0476074218750}, + {0x3dd0, 0.10156250}, + {0x3dd9, 0.105957031250}, + {0x3de9, 0.113769531250}, + {0x3df9, 0.121582031250}, + {0x3e1e, 0.1542968750}, + {0x3e77, 0.24121093750}, + {0x3e95, 0.2910156250}, + {0x3f38, 0.718750}, + {0x3fb3, 1.39843750}, + {0x3fc5, 1.53906250}, + {0x3fd3, 1.64843750}, + {0x3fd7, 1.67968750}, + {0x400b, 2.1718750}, + {0x40bf, 5.968750}, + {0x40c7, 6.218750}, + {0x4123, 10.18750}, + {0x412b, 10.68750}, + {0x41bf, 23.8750}, + {0x41ca, 25.250}, + {0x421b, 38.750}, + {0x4226, 41.50}, + {0x42a7, 83.50}, + {0x42b7, 91.50}, + {0x4311, 145.0}, + {0x431f, 159.0}, + {0x4334, 180.0}, + {0x434f, 207.0}, + {0x43b1, 354.0}, + {0x43e5, 458.0}, + {0x4476, 984.0}, + {0x4496, 1200.0}, + {0x44a4, 1312.0}, + {0x458b, 4448.0}, + {0x45a9, 5408.0}, + {0x45df, 7136.0}, + {0x45f6, 7872.0}, + {0x45fa, 8000.0}, + {0x4602, 8320.0}, + {0x4640, 12288.0}, + {0x4648, 12800.0}, + {0x46a7, 21376.0}, + {0x46b1, 22656.0}, + {0x4742, 49664.0}, + {0x4744, 50176.0}, + {0x475e, 56832.0}, + {0x477a, 64000.0}, + {0x4837, 187392.0}, + {0x488a, 282624.0}, + {0x488f, 292864.0}, + {0x48ea, 479232.0}, + {0x495c, 901120.0}, + {0x49aa, 1392640.0}, + {0x49b9, 1515520.0}, + {0x4a1e, 2588672.0}, + {0x4a2b, 2801664.0}, + {0x4a4c, 3342336.0}, + {0x4ab6, 5963776.0}, + {0x4b34, 11796480.0}, + {0x4b73, 15925248.0}, + {0x4b7b, 16449536.0}, + {0x4bcd, 26869760.0}, + {0x4bd0, 27262976.0}, + {0x4c07, 35389440.0}, + {0x4c17, 39583744.0}, + {0x4c53, 55312384.0}, + {0x4cad, 90701824.0}, + {0x4d1c, 163577856.0}, + {0x4dc0, 402653184.0}, + {0x4dde, 465567744.0}, + {0x4eef, 2004877312.0}, + {0x4efc, 2113929216.0}, + {0x4f12, 2449473536.0}, + {0x4f2f, 2936012800.0}, + {0x4f92, 4898947072.0}, + {0x4fad, 5804916736.0}, + {0x4fdc, 7381975040.0}, + {0x4feb, 7885291520.0}, + {0x5076, 16508780544.0}, + {0x5083, 17582522368.0}, + {0x5215, 159987531776.0}, + {0x52a9, 362924736512.0}, + {0x5394, 1271310319616.0}, + {0x53a0, 1374389534720.0}, + {0x53b7, 1571958030336.0}, + {0x540e, 2439541424128.0}, + {0x542f, 3006477107200.0}, + {0x5465, 3934190043136.0}, + {0x5529, 11613591568384.0}, + {0x554c, 14018773254144.0}, + {0x5596, 20615843020800.0}, + {0x55ae, 23914377904128.0}, + {0x55be, 26113401159680.0}, + {0x55da, 29961691856896.0}, + {0x568d, 77515569758208.0}, + {0x5690, 79164837199872.0}, + {0x56ad, 95107755802624.0}, + {0x5718, 167125767421952.0}, + {0x571c, 171523813933056.0}, + {0x571d, 172623325560832.0}, + {0x5826, 730075720843264.0}, + {0x587c, 1108307720798208.0}, + {0x5890, 1266637395197952.0}, + {0x58a8, 1477743627730944.0}, + {0x58f6, 2163838883463168.0}, + {0x5966, 4046202790215680.0}, + {0x5985, 4679521487814656.0}, + {0x59ad, 6086896371367936.0}, + {0x59b0, 6192449487634432.0}, + {0x59bc, 6614661952700416.0}, + {0x5a93, 20688410788233216.0}, + {0x5ab0, 24769797950537728.0}, + {0x5ab6, 25614222880669696.0}, + {0x5ae3, 31947409856659456.0}, + {0x5af5, 34480684647055360.0}, + {0x5afd, 35606584553897984.0}, + {0x5bb6, 102456891522678784.0}, + {0x5c3d, 212795082393255936.0}, + {0x5d57, 968273919884656640.0}, + {0x5d76, 1107885508333142016.0}, + {0x5d86, 1206964700135292928.0}, + {0x5db2, 1603281467343896576.0}, + {0x5dc1, 1738389456165011456.0}, + {0x5dc5, 1774418253183975424.0}, + {0x5e5d, 3981182070595518464.0}, + {0x5fa4, 23634890844440363008.0}, + {0x5fa8, 24211351596743786496.0}, + {0x5ff8, 35740566642812256256.0}, + {0x6006, 38622870404329373696.0}, + {0x6051, 60240148615707754496.0}, + {0x60ed, 136621198295911366656.0}, + {0x610b, 160256089140351729664.0}, + {0x6114, 170632382681813352448.0}, + {0x613b, 215596321361480384512.0}, + {0x6148, 230584300921369395200.0}, + {0x61ad, 398910840593969053696.0}, + {0x61fd, 583378281331064569856.0}, + {0x629f, 1466516153859909353472.0}, + {0x62a4, 1512633014044183232512.0}, + {0x62fb, 2315066381250548727808.0}, + {0x634b, 3744689046963038978048.0}, + {0x635a, 4021390208068682252288.0}, + {0x635f, 4113623928437230010368.0}, + {0x637c, 4648579506574807007232.0}, + {0x638d, 5201981828786093555712.0}, + {0x63a9, 6234999496913828446208.0}, + {0x6469, 17192365476697302106112.0}, + {0x64c0, 28334198897217871282176.0}, + {0x64d1, 30842956091242370301952.0}, + {0x64dd, 32613843522318487257088.0}, + {0x64ee, 35122600716342986276864.0}, + {0x64ef, 35270174668932662689792.0}, + {0x64fb, 37041062100008779644928.0}, + {0x6510, 42501298345826806923264.0}, + {0x6581, 76148159536273029070848.0}, + {0x65d4, 125142711796045598162944.0}, + {0x6612, 172366376624742050299904.0}, + {0x661c, 184172292831916163334144.0}, + {0x66c6, 467514281804094876155904.0}, + {0x66ed, 559600428220052957822976.0}, + {0x66f9, 587934627117270829105152.0}, + {0x6703, 618630009255923522994176.0}, + {0x6752, 991696961402625494876160.0}, + {0x6797, 1426154677826632854536192.0}, + {0x679c, 1473378342655329306673152.0}, + {0x67e3, 2143954383222818927017984.0}, + {0x6862, 4269019300514159273181184.0}, + {0x692f, 13222626152035006598348800.0}, + {0x693d, 14280436244197807126216704.0}, + {0x6943, 14733783426553293066731520.0}, + {0x695e, 16773845747152979799048192.0}, + {0x69ec, 35663311678631560653832192.0}, + {0x69f4, 36872237498246189828538368.0}, + {0x6a5c, 66490920078804604608839680.0}, + {0x6a7a, 75557863725914323419136000.0}, + {0x6ac3, 117870267412426344533852160.0}, + {0x6ad8, 130563988518379950868267008.0}, + {0x6ae3, 137213080526260411329150976.0}, + {0x6b37, 221233424989477138971230208.0}, + {0x6b77, 298604677444813406152425472.0}, + {0x6bd7, 519838102434290545123655680.0}, + {0x6be7, 558523728661958678714253312.0}, + {0x6c15, 720519788490318988124880896.0}, + {0x6c33, 865590886844074489089622016.0}, + {0x6c43, 942962139299410756270817280.0}, + {0x6c45, 952633545856327789668466688.0}, + {0x6c48, 967140655691703339764940800.0}, + {0x6da6, 6421813953792910176039206912.0}, + {0x6dba, 7195526478346272847851159552.0}, + {0x6def, 9245864668412683928152834048.0}, + {0x6e24, 12688885402675147817716023296.0}, + {0x6e28, 12998370412496492886440804352.0}, + {0x6e3d, 14623166714058554497245904896.0}, + {0x6ea0, 24758800785707605497982484480.0}, + {0x6eef, 36983458673650735712611336192.0}, + {0x6ef9, 38530883722757461056235241472.0}, + {0x6f55, 65920307091946499638378364928.0}, + {0x6f5c, 68086702160695915119451832320.0}, + {0x6f65, 70872067249088020737974861824.0}, + {0x6f9e, 97797263103545041717030813696.0}, + {0x6fdc, 136173404321391830238903664640.0}, + {0x70ab, 423375493435600054015500484608.0}, + {0x70dc, 544693617285567320955614658560.0}, + {0x714a, 1000255551742587262118492372992.0}, + {0x71ba, 1842054778456645849049896845312.0}, + {0x71d3, 2089642786313721904029721690112.0}, + {0x71ee, 2357037834799364043407932522496.0}, + {0x71f6, 2436265997313628381001476472832.0}, + {0x7251, 4139671491370311639262671405056.0}, + {0x72b8, 7288990951312319058606043430912.0}, + {0x731f, 12597277839768029677373488103424.0}, + {0x7328, 13310331302396408715715383656448.0}, + {0x7356, 16954826778052568245018405371904.0}, + {0x7358, 17113283103081096920205493272576.0}, + {0x7375, 19410899815994762710418267832320.0}, + {0x737c, 19965496953594613073573075484672.0}, + {0x7467, 73206822163180247936434610110464.0}, + {0x74a4, 103947349218714810922729662840832.0}, + {0x74b1, 112187078120198302032458233675776.0}, + {0x7530, 223106505640168374663419764146176.0}, + {0x7563, 287756686251808074139751627620352.0}, + {0x756f, 302968493454546826957712066084864.0}, + {0x7571, 305503794655003285760705472495616.0}, + {0x7626, 841719998551544322593810928369664.0}, + {0x7660, 1135814937804493543741046072016896.0}, + {0x7675, 1242297588223664813466769141268480.0}, + {0x76b2, 1805134454724998667731305364455424.0}, + {0x772c, 3488574451828087312918927221194752.0}, + {0x77b2, 7220537818899994670925221457821696.0}, + {0x783a, 15090112745116842795416754956795904.0}, + {0x795d, 71718600358512306619077480547352576.0}, + {0x796e, 77235415770705560974391132897148928.0}, + {0x798c, 90865195024359483499283685761351680.0}, + {0x79af, 113581493780449354374104607201689600.0}, + {0x7aa5, 428364490829123279353765947160657920.0}, + {0x7acf, 537402724858354659552906370074279936.0}, + {0x7adc, 571152654438831039138354596214210560.0}, + {0x7b07, 700960075902201729851617004444712960.0}, + {0x7b0f, 742498450770480350879860975078473728.0}, + {0x7b12, 758075341346084833765452464066134016.0}, + {0x7b30, 913844247102129662621367353942736896.0}, + {0x7bc2, 2014611181111513119869832575737397248.0}, + {0x7bd1, 2170380086867557948725747465614000128.0}, + {0x7bd5, 2211918461735836569753991436247760896.0}, + {0x7cf1, 10010748343255147667806796922736345088.0}, + {0x7d2f, 14538431203897517359885389721816268800.0}, + {0x7da0, 26584559915698317458076141205606891520.0}, + {0x7e58, 71778311772385457136805581255138607104.0}, + {0x7e81, 85735205728127073802295555388082225152.0}, + {0x7f09, 182104235422533474587821567258407206912.0}, + {0x7f24, 217993391308726203156224357885976510464.0}, + {0x7f86, std::numeric_limits::quiet_NaN()}, + {0x7f88, std::numeric_limits::quiet_NaN()}, + {0x7f8f, std::numeric_limits::quiet_NaN()}, + {0x7fa0, std::numeric_limits::quiet_NaN()}, + {0x7fcd, std::numeric_limits::quiet_NaN()}, + {0x8023, -0.00000000000000000000000000000000000000321424236553}, + {0x8074, -0.00000000000000000000000000000000000001065291755433}, + {0x8080, -0.00000000000000000000000000000000000001175494350822}, + {0x80a5, -0.00000000000000000000000000000000000001515285686607}, + {0x80d2, -0.00000000000000000000000000000000000001928545419318}, + {0x80fd, -0.00000000000000000000000000000000000002323438052797}, + {0x810a, -0.00000000000000000000000000000000000002534659693961}, + {0x8124, -0.00000000000000000000000000000000000003012204273982}, + {0x81e3, -0.00000000000000000000000000000000000008338663051146}, + {0x81f1, -0.00000000000000000000000000000000000008852941829630}, + {0x8285, -0.00000000000000000000000000000000000019542593582421}, + {0x828b, -0.00000000000000000000000000000000000020424214345537}, + {0x829f, -0.00000000000000000000000000000000000023362950222593}, + {0x82bc, -0.00000000000000000000000000000000000027624117244324}, + {0x82ee, -0.00000000000000000000000000000000000034970956936963}, + {0x82f9, -0.00000000000000000000000000000000000036587261669344}, + {0x8393, -0.00000000000000000000000000000000000086398834785438}, + {0x8394, -0.00000000000000000000000000000000000086986581960849}, + {0x843e, -0.00000000000000000000000000000000000223343926656235}, + {0x8451, -0.00000000000000000000000000000000000245678319321858}, + {0x847f, -0.00000000000000000000000000000000000299751059459683}, + {0x8483, -0.00000000000000000000000000000000000307979519915439}, + {0x84a0, -0.00000000000000000000000000000000000376158192263132}, + {0x84aa, -0.00000000000000000000000000000000000399668079279578}, + {0x84ba, -0.00000000000000000000000000000000000437283898505891}, + {0x84e4, -0.00000000000000000000000000000000000536025423974963}, + {0x8510, -0.00000000000000000000000000000000000677084746073638}, + {0x854d, -0.00000000000000000000000000000000000963905367674276}, + {0x8557, -0.00000000000000000000000000000000001010925141707167}, + {0x8584, -0.00000000000000000000000000000000001241322034468336}, + {0x85b7, -0.00000000000000000000000000000000001720923729603829}, + {0x8656, -0.00000000000000000000000000000000004024892657215512}, + {0x87c7, -0.00000000000000000000000000000000029942192104145307}, + {0x88d9, -0.00000000000000000000000000000000130602124353759431}, + {0x896c, -0.00000000000000000000000000000000284074666797117288}, + {0x89a0, -0.00000000000000000000000000000000385185988877447171}, + {0x89b2, -0.00000000000000000000000000000000428519412626159977}, + {0x89fe, -0.00000000000000000000000000000000611482757342947383}, + {0x8a31, -0.00000000000000000000000000000000852224000391351865}, + {0x8a55, -0.00000000000000000000000000000001025557695386203092}, + {0x8a68, -0.00000000000000000000000000000001117039367744596795}, + {0x8b4a, -0.00000000000000000000000000000003890378487662216423}, + {0x8b4c, -0.00000000000000000000000000000003928897086549961140}, + {0x8b5b, -0.00000000000000000000000000000004217786578208046518}, + {0x8bbf, -0.00000000000000000000000000000007357052387559240959}, + {0x8bff, -0.00000000000000000000000000000009822242716374902851}, + {0x8c43, -0.00000000000000000000000000000015022253566220439654}, + {0x8c6b, -0.00000000000000000000000000000018103741477240017019}, + {0x8d1d, -0.00000000000000000000000000000048379360203007364629}, + {0x8d23, -0.00000000000000000000000000000050228252949619111048}, + {0x8d30, -0.00000000000000000000000000000054234187233944561622}, + {0x8dee, -0.00000000000000000000000000000146678824564531882569}, + {0x8e54, -0.00000000000000000000000000000261310174854460160543}, + {0x8e68, -0.00000000000000000000000000000285962078142616779462}, + {0x8eb0, -0.00000000000000000000000000000433873497871556492976}, + {0x8efe, -0.00000000000000000000000000000626158343519178120546}, + {0x8f3d, -0.00000000000000000000000000000931841944292320195143}, + {0x8f95, -0.00000000000000000000000000001469253435974134487579}, + {0x8fa3, -0.00000000000000000000000000001607304094387811553526}, + {0x8fbe, -0.00000000000000000000000000001873544649899903037853}, + {0x9009, -0.00000000000000000000000000002701848600381965433535}, + {0x9017, -0.00000000000000000000000000002977949917209319565429}, + {0x90a8, -0.00000000000000000000000000006626431603856499165459}, + {0x90d2, -0.00000000000000000000000000008283039504820623956823}, + {0x9260, -0.00000000000000000000000000070681937107802657764891}, + {0x92d0, -0.00000000000000000000000000131266454628776364420512}, + {0x92ed, -0.00000000000000000000000000149568027629903838306064}, + {0x9356, -0.00000000000000000000000000270105973947674442172976}, + {0x94af, -0.00000000000000000000000001767048427695066444122272}, + {0x9503, -0.00000000000000000000000002645523931749185190628773}, + {0x953c, -0.00000000000000000000000003796629764647685617085567}, + {0x9556, -0.00000000000000000000000004321695583162791074767614}, + {0x9598, -0.00000000000000000000000006139231108792002274436236}, + {0x9599, -0.00000000000000000000000006179620787139318078873317}, + {0x95cf, -0.00000000000000000000000008360663417894371518475664}, + {0x95df, -0.00000000000000000000000009006898271451424389468952}, + {0x95f4, -0.00000000000000000000000009855081516745056282647643}, + {0x960f, -0.00000000000000000000000011551448007332320069005024}, + {0x9615, -0.00000000000000000000000012036124147500109722249990}, + {0x966a, -0.00000000000000000000000018902369466543796476553675}, + {0x968c, -0.00000000000000000000000022618219874496850484765081}, + {0x96b3, -0.00000000000000000000000028919009696678115976949640}, + {0x96d7, -0.00000000000000000000000034735123378691591815889232}, + {0x9719, -0.00000000000000000000000049436966297114544630986535}, + {0x9761, -0.00000000000000000000000072701421025168447986744905}, + {0x97f0, -0.00000000000000000000000155096364853692689038389130}, + {0x97fd, -0.00000000000000000000000163497417949934376361301874}, + {0x983c, -0.00000000000000000000000242984304937451879493476303}, + {0x987d, -0.00000000000000000000000326994835899868752722603749}, + {0x9895, -0.00000000000000000000000385155972720003511111999672}, + {0x98b3, -0.00000000000000000000000462704155146849855631194237}, + {0x98fa, -0.00000000000000000000000646234853557052870993288041}, + {0x998f, -0.00000000000000000000001478585344938536968832643037}, + {0x999b, -0.00000000000000000000001602662436821491120063354341}, + {0x99e4, -0.00000000000000000000002357464745776128873383514772}, + {0x9a8b, -0.00000000000000000000005748905257243542340356290410}, + {0x9a95, -0.00000000000000000000006162495563520056177791994756}, + {0x9a9f, -0.00000000000000000000006576085869796570015227699102}, + {0x9bb2, -0.00000000000000000000029447629806887785225422149438}, + {0x9bc4, -0.00000000000000000000032425480012078684854959220729}, + {0x9c09, -0.00000000000000000000045329497567905916582953196325}, + {0x9c79, -0.00000000000000000000082387189010281556417192305730}, + {0x9cca, -0.00000000000000000000133672386988569272259219644639}, + {0x9d21, -0.00000000000000000000213081725793659929046874879077}, + {0x9d2d, -0.00000000000000000000228963593554678060404405925965}, + {0x9d53, -0.00000000000000000000279256174797902143036587574443}, + {0x9d59, -0.00000000000000000000287197108678411208715353097887}, + {0x9d9b, -0.00000000000000000000410281583826301726736218711267}, + {0x9dc9, -0.00000000000000000000532042569994107400477290070739}, + {0x9e3a, -0.00000000000000000000984675801183124144166924907040}, + {0x9eb4, -0.00000000000000000001905824131322175762903725626529}, + {0x9eb7, -0.00000000000000000001937587866844212025618787720305}, + {0x9ec3, -0.00000000000000000002064642808932357076479036095407}, + {0x9ee0, -0.00000000000000000002371692252312040949391303001903}, + {0x9f0a, -0.00000000000000000002922263668027336169785712627345}, + {0x9f5b, -0.00000000000000000004637505386217294356399065691221}, + {0x9fb7, -0.00000000000000000007750351467376848102475150881219}, + {0xa007, -0.00000000000000000011434944787933054577422353759175}, + {0xa06c, -0.00000000000000000019989977555201488002012411016040}, + {0xa07a, -0.00000000000000000021175823681357508476708062516991}, + {0xa0c0, -0.00000000000000000032526065174565133020223584026098}, + {0xa0e5, -0.00000000000000000038794108984246955529329170531128}, + {0xa108, -0.00000000000000000046078592330633938445316744036973}, + {0xa128, -0.00000000000000000056920614055488982785391272045672}, + {0xa178, -0.00000000000000000084025668367626593635577592067420}, + {0xa1f6, -0.00000000000000000166696084019646306728645868133754}, + {0xa305, -0.00000000000000000720994444702860448614956112578511}, + {0xa312, -0.00000000000000000791467585914418236825440544635057}, + {0xa3f2, -0.00000000000000002623769257414920730298035778105259}, + {0xa3f3, -0.00000000000000002634611279139775774638110306113958}, + {0xa40a, -0.00000000000000002992397996059992237860569730401039}, + {0xa412, -0.00000000000000003165870343657672947301762178540230}, + {0xa413, -0.00000000000000003187554387107383035981911234557629}, + {0xa421, -0.00000000000000003491130995403324277503998018801212}, + {0xa450, -0.00000000000000004510281037539698445471003651618958}, + {0xa456, -0.00000000000000004640385298237958977551897987723351}, + {0xa4a4, -0.00000000000000007112366251504909087088890373706818}, + {0xa4c9, -0.00000000000000008716985466783455649419920518994331}, + {0xa4ed, -0.00000000000000010278236595162582034390652552247047}, + {0xa522, -0.00000000000000014051260155412137464736588299274445}, + {0xa53a, -0.00000000000000016132928326584305978030897676944733}, + {0xa541, -0.0000000000000001674008154317618846107507124543190}, + {0xa6f7, -0.00000000000000171390679426508540927898138761520386}, + {0xa76c, -0.00000000000000327515792264421179424971342086791992}, + {0xa7a3, -0.00000000000000452415882534751290222629904747009277}, + {0xa7ae, -0.00000000000000482947015711943095084279775619506836}, + {0xa7bc, -0.00000000000000521804821573823573999106884002685547}, + {0xa837, -0.00000000000001015854067532018234487622976303100586}, + {0xa83e, -0.00000000000001054711873393898713402450084686279297}, + {0xa881, -0.00000000000001432187701766451937146484851837158203}, + {0xa8b9, -0.00000000000002053912595556539599783718585968017578}, + {0xa8c2, -0.00000000000002153832667772803688421845436096191406}, + {0xa8ca, -0.00000000000002242650509742816211655735969543457031}, + {0xa8d6, -0.00000000000002375877272697834996506571769714355469}, + {0xa8e4, -0.00000000000002531308496145356912165880203247070312}, + {0xa92d, -0.00000000000003841371665203041629865765571594238281}, + {0xa933, -0.00000000000003974598428158060414716601371765136719}, + {0xa937, -0.00000000000004063416270128072937950491905212402344}, + {0xa950, -0.0000000000000461852778244065120816230773925781250}, + {0xa956, -0.00000000000004751754545395669993013143539428710938}, + {0xa9b8, -0.0000000000000817124146124115213751792907714843750}, + {0xa9c6, -0.00000000000008792966355031239800155162811279296875}, + {0xa9f0, -0.000000000000106581410364015027880668640136718750}, + {0xaa2d, -0.00000000000015365486660812166519463062286376953125}, + {0xaae6, -0.0000000000004085620730620576068758964538574218750}, + {0xaaff, -0.00000000000045297099404706386849284172058105468750}, + {0xab12, -0.000000000000518696197104873135685920715332031250}, + {0xab42, -0.000000000000689226453687297180294990539550781250}, + {0xabfa, -0.00000000000177635683940025046467781066894531250}, + {0xac3a, -0.0000000000026432189770275726914405822753906250}, + {0xadbb, -0.00000000002125943865394219756126403808593750}, + {0xadc3, -0.00000000002216893335571512579917907714843750}, + {0xadda, -0.0000000000247837306233122944831848144531250}, + {0xaddf, -0.00000000002535216481192037463188171386718750}, + {0xae16, -0.000000000034106051316484808921813964843750}, + {0xae77, -0.0000000000561612978344783186912536621093750}, + {0xaee2, -0.00000000010277290130034089088439941406250}, + {0xb03e, -0.00000000069121597334742546081542968750}, + {0xb050, -0.00000000075669959187507629394531250}, + {0xb075, -0.000000000891304807737469673156738281250}, + {0xb11d, -0.0000000022846506908535957336425781250}, + {0xb125, -0.0000000024010660126805305480957031250}, + {0xb139, -0.0000000026921043172478675842285156250}, + {0xb155, -0.0000000030995579436421394348144531250}, + {0xb18d, -0.000000004103640094399452209472656250}, + {0xb23c, -0.000000010943040251731872558593750}, + {0xb2a8, -0.0000000195577740669250488281250}, + {0xb341, -0.000000044936314225196838378906250}, + {0xb369, -0.000000054249539971351623535156250}, + {0xb37b, -0.000000058440491557121276855468750}, + {0xb3c6, -0.0000000922009348869323730468750}, + {0xb3c9, -0.00000009359791874885559082031250}, + {0xb3dc, -0.000000102445483207702636718750}, + {0xb3e2, -0.0000001052394509315490722656250}, + {0xb404, -0.00000012293457984924316406250}, + {0xb42d, -0.0000001611188054084777832031250}, + {0xb487, -0.000000251457095146179199218750}, + {0xb499, -0.000000284984707832336425781250}, + {0xb49b, -0.000000288709998130798339843750}, + {0xb4be, -0.00000035390257835388183593750}, + {0xb599, -0.0000011399388313293457031250}, + {0xb5be, -0.000001415610313415527343750}, + {0xb661, -0.000003352761268615722656250}, + {0xb67f, -0.000003799796104431152343750}, + {0xb6f4, -0.000007271766662597656250}, + {0xb70f, -0.0000085234642028808593750}, + {0xb729, -0.0000100731849670410156250}, + {0xb731, -0.0000105500221252441406250}, + {0xb735, -0.0000107884407043457031250}, + {0xb76f, -0.0000142455101013183593750}, + {0xb770, -0.000014305114746093750}, + {0xb7a4, -0.0000195503234863281250}, + {0xb7b1, -0.000021100044250488281250}, + {0xb829, -0.00004029273986816406250}, + {0xb882, -0.000061988830566406250}, + {0xb9a6, -0.0003166198730468750}, + {0xb9c5, -0.00037574768066406250}, + {0xb9cc, -0.000389099121093750}, + {0xb9d3, -0.00040245056152343750}, + {0xb9dd, -0.00042152404785156250}, + {0xbb04, -0.002014160156250}, + {0xbb14, -0.002258300781250}, + {0xbb19, -0.00233459472656250}, + {0xbb33, -0.00273132324218750}, + {0xbb66, -0.0035095214843750}, + {0xbbc5, -0.0060119628906250}, + {0xbc0d, -0.008605957031250}, + {0xbcb0, -0.0214843750}, + {0xbcc8, -0.02441406250}, + {0xbce0, -0.027343750}, + {0xbce8, -0.02832031250}, + {0xbd06, -0.032714843750}, + {0xbd77, -0.0603027343750}, + {0xbe31, -0.17285156250}, + {0xbe3a, -0.1816406250}, + {0xbe5d, -0.21582031250}, + {0xbe85, -0.2597656250}, + {0xbe9a, -0.300781250}, + {0xbea5, -0.3222656250}, + {0xbeb0, -0.343750}, + {0xbebf, -0.3730468750}, + {0xbeee, -0.464843750}, + {0xbf2b, -0.667968750}, + {0xbfac, -1.343750}, + {0xc022, -2.531250}, + {0xc026, -2.593750}, + {0xc05e, -3.468750}, + {0xc07e, -3.968750}, + {0xc07f, -3.9843750}, + {0xc086, -4.18750}, + {0xc0ae, -5.43750}, + {0xc0c2, -6.06250}, + {0xc0e6, -7.18750}, + {0xc13e, -11.8750}, + {0xc198, -19.0}, + {0xc1be, -23.750}, + {0xc1c1, -24.1250}, + {0xc1eb, -29.3750}, + {0xc225, -41.250}, + {0xc276, -61.50}, + {0xc27f, -63.750}, + {0xc29f, -79.50}, + {0xc313, -147.0}, + {0xc31b, -155.0}, + {0xc324, -164.0}, + {0xc35b, -219.0}, + {0xc394, -296.0}, + {0xc39d, -314.0}, + {0xc3b5, -362.0}, + {0xc3be, -380.0}, + {0xc429, -676.0}, + {0xc444, -784.0}, + {0xc44b, -812.0}, + {0xc4b5, -1448.0}, + {0xc4eb, -1880.0}, + {0xc523, -2608.0}, + {0xc557, -3440.0}, + {0xc55e, -3552.0}, + {0xc56d, -3792.0}, + {0xc58b, -4448.0}, + {0xc64d, -13120.0}, + {0xc6b8, -23552.0}, + {0xc6ca, -25856.0}, + {0xc777, -63232.0}, + {0xc7d6, -109568.0}, + {0xc868, -237568.0}, + {0xc8ca, -413696.0}, + {0xc910, -589824.0}, + {0xc9c5, -1613824.0}, + {0xc9c8, -1638400.0}, + {0xc9df, -1826816.0}, + {0xca3a, -3047424.0}, + {0xca42, -3178496.0}, + {0xca6b, -3850240.0}, + {0xcaa0, -5242880.0}, + {0xcaa2, -5308416.0}, + {0xcaac, -5636096.0}, + {0xcb3a, -12189696.0}, + {0xcb84, -17301504.0}, + {0xcc50, -54525952.0}, + {0xcc89, -71827456.0}, + {0xcc94, -77594624.0}, + {0xccaf, -91750400.0}, + {0xcce0, -117440512.0}, + {0xcce1, -117964800.0}, + {0xcd6d, -248512512.0}, + {0xcda8, -352321536.0}, + {0xcdba, -390070272.0}, + {0xcdd0, -436207616.0}, + {0xcde5, -480247808.0}, + {0xcdf7, -517996544.0}, + {0xce30, -738197504.0}, + {0xcec2, -1627389952.0}, + {0xcf03, -2197815296.0}, + {0xcf25, -2768240640.0}, + {0xcf57, -3607101440.0}, + {0xd036, -12213813248.0}, + {0xd09e, -21206401024.0}, + {0xd103, -35165044736.0}, + {0xd104, -35433480192.0}, + {0xd11f, -42681237504.0}, + {0xd125, -44291850240.0}, + {0xd19c, -83751862272.0}, + {0xd1c7, -106837311488.0}, + {0xd1cf, -111132278784.0}, + {0xd1d8, -115964116992.0}, + {0xd231, -190052302848.0}, + {0xd28a, -296352743424.0}, + {0xd294, -317827579904.0}, + {0xd2be, -408021893120.0}, + {0xd2c1, -414464344064.0}, + {0xd2c6, -425201762304.0}, + {0xd2db, -470298918912.0}, + {0xd334, -773094113280.0}, + {0xd36f, -1026497183744.0}, + {0xd375, -1052266987520.0}, + {0xd3c3, -1675037245440.0}, + {0xd3d5, -1829656068096.0}, + {0xd3f2, -2078764171264.0}, + {0xd44c, -3504693313536.0}, + {0xd49b, -5325759447040.0}, + {0xd4cd, -7043746365440.0}, + {0xd4e8, -7971459301376.0}, + {0xd538, -12644383719424.0}, + {0xd54c, -14018773254144.0}, + {0xd554, -14568529068032.0}, + {0xd5a3, -22402549415936.0}, + {0xd5bf, -26250840113152.0}, + {0xd64a, -55525337202688.0}, + {0xd74f, -227598906949632.0}, + {0xd75f, -245191092994048.0}, + {0xd762, -248489627877376.0}, + {0xd7d6, -470590976688128.0}, + {0xd7db, -481586092965888.0}, + {0xd819, -672901116198912.0}, + {0xd82d, -760862046420992.0}, + {0xd85b, -963172185931776.0}, + {0xd936, -3201777860083712.0}, + {0xd967, -4063794976260096.0}, + {0xd976, -4327677766926336.0}, + {0xd985, -4679521487814656.0}, + {0xd98c, -4925812092436480.0}, + {0xd9c9, -7072058789855232.0}, + {0xda1b, -10907155347537920.0}, + {0xda5a, -15340386230730752.0}, + {0xdab6, -25614222880669696.0}, + {0xdacc, -28710447624486912.0}, + {0xdb00, -36028797018963968.0}, + {0xdb0d, -39687971716202496.0}, + {0xdb24, -46161896180547584.0}, + {0xdb4c, -57420895248973824.0}, + {0xdbbd, -106397541196627968.0}, + {0xdbbf, -107523441103470592.0}, + {0xdbcc, -114841790497947648.0}, + {0xdbf4, -137359788634800128.0}, + {0xdc88, -306244774661193728.0}, + {0xdc9c, -351280770934898688.0}, + {0xdd24, -738590338888761344.0}, + {0xdd71, -1085367510196289536.0}, + {0xdd86, -1206964700135292928.0}, + {0xdd8f, -1288029493427961856.0}, + {0xdea9, -6088866696204910592.0}, + {0xded6, -7710162562058289152.0}, + {0xdedf, -8034421735228964864.0}, + {0xdef6, -8863084066665136128.0}, + {0xdf62, -16285016252571713536.0}, + {0xdf6a, -16861477004875137024.0}, + {0xdf71, -17365880163140632576.0}, + {0xdfc9, -28967152803247030272.0}, + {0xdff0, -34587645138205409280.0}, + {0xdff5, -35308221078584688640.0}, + {0xe037, -52746158835763249152.0}, + {0xe07e, -73210515542534782976.0}, + {0xe09a, -88774955854727217152.0}, + {0xe09d, -90504338111637487616.0}, + {0xe0cc, -117597993469898391552.0}, + {0xe113, -169479461177206505472.0}, + {0xe18f, -329735550317558235136.0}, + {0xe1aa, -391993311566327971840.0}, + {0xe21a, -710199646837817737216.0}, + {0xe233, -825491797298502434816.0}, + {0xe238, -848550227390639374336.0}, + {0xe247, -917725517667050192896.0}, + {0xe262, -1042241040164589666304.0}, + {0xe26a, -1079134528312008769536.0}, + {0xe271, -1111416330441000484864.0}, + {0xe28b, -1282048713122813837312.0}, + {0xe290, -1328165573307087716352.0}, + {0xe2b2, -1641760222560150093824.0}, + {0xe2eb, -2167492428660872314880.0}, + {0xe2f5, -2259726149029420072960.0}, + {0xe34b, -3744689046963038978048.0}, + {0xe363, -4187410904732068216832.0}, + {0xe388, -5017514388048998039552.0}, + {0xe3a0, -5902958103587056517120.0}, + {0xe3bd, -6972869259862210510848.0}, + {0xe3d6, -7895206463547688091648.0}, + {0xe3e4, -8411715297611555536896.0}, + {0xe406, -9887454823508319666176.0}, + {0xe5a3, -96218217088469021229056.0}, + {0xe5ae, -102711471002414783397888.0}, + {0xe5d7, -126913599227121715118080.0}, + {0xe5d9, -128094190847839126421504.0}, + {0xe644, -231395957660612615471104.0}, + {0xe66e, -280980805730743890214912.0}, + {0xe6b8, -434457716424007359660032.0}, + {0xe7f4, -2304514843640386864283648.0}, + {0xe824, -3097872412762487260184576.0}, + {0xe827, -3154540810556923002748928.0}, + {0xe880, -4835703278458516698824704.0}, + {0xe8f2, -9142501510835633133715456.0}, + {0xe928, -12693721105953606334414848.0}, + {0xe980, -19342813113834066795298816.0}, + {0xe9d1, -31583187037432187189198848.0}, + {0xe9fa, -37778931862957161709568000.0}, + {0xea15, -45032486780644936757805056.0}, + {0xea77, -74651169361203351538106368.0}, + {0xeaca, -122101507781077546645323776.0}, + {0xeacd, -123914896510499490407383040.0}, + {0xead9, -131168451428187265455620096.0}, + {0xeada, -131772914337994580042973184.0}, + {0xeb79, -301022529084042664501837824.0}, + {0xeba3, -394109817194369110954213376.0}, + {0xebf3, -587537948332709778907201536.0}, + {0xec09, -662491349148816787738984448.0}, + {0xec13, -710848381933401954727231488.0}, + {0xec48, -967140655691703339764940800.0}, + {0xece9, -2253437727761668781652312064.0}, + {0xed1f, -3075507285099616620452511744.0}, + {0xed24, -3172221350668786954429005824.0}, + {0xed9f, -6151014570199233240905023488.0}, + {0xede6, -8897694032363670725837455360.0}, + {0xedfe, -9826149061827705932011798528.0}, + {0xeea3, -25223028300439623101069656064.0}, + {0xeea5, -25532513310260968169794437120.0}, + {0xeec4, -30329530962491816735028543488.0}, + {0xeee0, -34662321099990647697175478272.0}, + {0xef23, -50446056600879246202139312128.0}, + {0xefa0, -99035203142830421991929937920.0}, + {0xefab, -105843873358900013503875121152.0}, + {0xf020, -198070406285660843983859875840.0}, + {0xf099, -378809652021326364119132012544.0}, + {0xf0a9, -418423733278458532915903987712.0}, + {0xf0cc, -505079536028435152158842683392.0}, + {0xf0e6, -569452418071274926453597143040.0}, + {0xf16b, -1163663636928257458405176770560.0}, + {0xf19b, -1535045648713871540874914037760.0}, + {0xf1bc, -1861861819085211933448282832896.0}, + {0xf1da, -2158967428513703199424072646656.0}, + {0xf231, -3505846191256196938514319802368.0}, + {0xf253, -4179285572627443808059443380224.0}, + {0xf287, -5347900969712842787564216647680.0}, + {0xf296, -5942112188569825319515796275200.0}, + {0xf2f3, -9626221745483117017615589965824.0}, + {0xf339, -14657210065138902454805630812160.0}, + {0xf36e, -18856302678394912347263460179968.0}, + {0xf3b3, -28363682180106632858488734220288.0}, + {0xf3c0, -30423614405477505635920876929024.0}, + {0xf3da, -34543478856219251190785162346496.0}, + {0xf3eb, -37237236381704238668965656657920.0}, + {0xf407, -42783207757702742300513733181440.0}, + {0xf418, -48170722808672717256874721804288.0}, + {0xf425, -52290587259414462811739007221760.0}, + {0xf473, -77009773963864936140924719726592.0}, + {0xf515, -188879939434006180823008777601024.0}, + {0xf57a, -316912650057057350374175801344000.0}, + {0xf594, -375224577667555902843024148791296.0}, + {0xf61c, -791013974542415146533942800154624.0}, + {0xf634, -912708432164325169077626307870720.0}, + {0xf69f, -1612451563490307798703806477238272.0}, + {0xf6b2, -1805134454724998667731305364455424.0}, + {0xf6f0, -2433889152438200450873670154321920.0}, + {0xf71e, -3204620717376963926983665703190528.0}, + {0xf735, -3671116138260952346734452482768896.0}, + {0xf748, -4056481920730334084789450257203200.0}, + {0xf75f, -4522977341614322504540237036781568.0}, + {0xf796, -6084722881095501127184175385804800.0}, + {0xf8a3, -26448262123161778232827215676964864.0}, + {0xf8d0, -33749929580476379585448226139930624.0}, + {0xf8ef, -38779967162181993850587144458862592.0}, + {0xf9a3, -105793048492647112931308862707859456.0}, + {0xfa0d, -183028464263352673905699995605008384.0}, + {0xfa18, -197307280624323449884158860510363648.0}, + {0xfa31, -229759135990166122562474462567989248.0}, + {0xfaa9, -438749084546192934610826939819098112.0}, + {0xfab5, -469902865697401900382009917794418688.0}, + {0xfacd, -532210427999819831924375873745059840.0}, + {0xfb84, -1370766370653194493932051030914105344.0}, + {0xfba8, -1744611744467702083186246766617952256.0}, + {0xfbfa, -2596148429267413814265248164610048000.0}, + {0xfc20, -3323069989462289682259517650700861440.0}, + {0xfc76, -5109220108798270386474008387952574464.0}, + {0xfc93, -6106141105636957291151863683162832896.0}, + {0xfccf, -8598443597733674552846501921188478976.0}, + {0xfcdc, -9138442471021296626213673539427368960.0}, + {0xfcf0, -9969209968386869046778552952102584320.0}, + {0xfd5b, -18193808192306036010370859137587216384.0}, + {0xfda9, -28079941410956347815092924148422279168.0}, + {0xfdca, -33563006893569125790821128272078700544.0}, + {0xfe1d, -52172198834557948011474427116003524608.0}, + {0xfe24, -54498347827181550789056089471494127616.0}, + {0xfe4b, -67458320786084480549868208309227487232.0}, + {0xfe4f, -68787548781869396422772015369507831808.0}, + {0xfe64, -75765995759740204755517002435979640832.0}, + {0xfea2, -107667467658578185705208371882707910656.0}, + {0xfec2, -128935115591136839671669284847193423872.0}, + {0xfee3, -150867377521587951574582101341819109376.0}, + {0xfee9, -154855061508942699193293522522660143104.0}, + {0xff57, -285784019093756912674318517960274083840.0}, + {0xff60, -297747071055821155530452781502797185024.0}, + {0xff8e, std::numeric_limits::quiet_NaN()}, + {0xffb0, std::numeric_limits::quiet_NaN()}, + {0xfffa, std::numeric_limits::quiet_NaN()}, + }; + return result; +} + +TEST_CASE(check_bf16_values) +{ + for(auto [x, f] : bf16_lut()) + { + + auto h = migraphx::bit_cast(x); + if(std::isnan(f)) + { + CHECK(std::isnan(h)); + } + else if(std::isinf(f)) + { + CHECK(std::isinf(h)); + CHECK((h < 0) == (f < 0)); + CHECK(bit_equal(x, migraphx::bf16(f))); + } + else + { + CHECK(bit_equal(x, migraphx::bf16(f))); + CHECK(migraphx::float_equal(float(h), f)); + } + } +} + +TEST_CASE(check_flows) +{ + // check positive underflow + CHECK(bit_equal(std::numeric_limits::min() * + std::numeric_limits::min(), + migraphx::bf16(0))); + + // check overflow + CHECK(bit_equal(std::numeric_limits::infinity() + + std::numeric_limits::infinity(), + std::numeric_limits::infinity())); + CHECK(bit_equal(std::numeric_limits::max() + + std::numeric_limits::max(), + std::numeric_limits::infinity())); + CHECK(bit_equal(std::numeric_limits::max() / + std::numeric_limits::epsilon(), + std::numeric_limits::infinity())); + CHECK(bit_equal(std::numeric_limits::max() + + std::numeric_limits::min(), + std::numeric_limits::max())); + + // check negative underflow + CHECK(bit_equal(std::numeric_limits::lowest() + + std::numeric_limits::lowest(), + -std::numeric_limits::infinity())); + CHECK(bit_equal(-std::numeric_limits::infinity() - + std::numeric_limits::infinity(), + -std::numeric_limits::infinity())); + CHECK(bit_equal(std::numeric_limits::lowest() - + std::numeric_limits::min(), + std::numeric_limits::lowest())); +} + +TEST_CASE(test_nan) +{ + float f_qnan = std::numeric_limits::quiet_NaN(); + migraphx::bf16 bf16_qnan(f_qnan); + EXPECT(bf16_qnan.is_nan()); + EXPECT(std::isnan(bf16_qnan)); + + float f_snan = std::numeric_limits::signaling_NaN(); + migraphx::bf16 bf16_snan(f_snan); + EXPECT(bf16_snan.is_nan()); + EXPECT(std::isnan(bf16_snan)); +} + +TEST_CASE(test_bool) +{ + float zero = 0.0; + float two = 2.0; + float other = -0.375; + migraphx::bf16 bf16_zero(zero); + migraphx::bf16 bf16_two(two); + migraphx::bf16 bf16_other(other); + EXPECT(not static_cast(bf16_zero)); + EXPECT(static_cast(bf16_two)); + EXPECT(static_cast(bf16_other)); +} + +TEST_CASE(test_pos_infinity) +{ + float finf = std::numeric_limits::infinity(); + migraphx::bf16 bf16_inf_1(finf); + CHECK(bit_equal(bf16_inf_1, std::numeric_limits::infinity())); +} + +TEST_CASE(test_neg_infinity) +{ + float finf = -1.0 * std::numeric_limits::infinity(); + migraphx::bf16 bf16_neginf_1(finf); + CHECK(bit_equal(bf16_neginf_1, -std::numeric_limits::infinity())); +} + +TEST_CASE(test_numeric_max_1) +{ + float fmax = std::numeric_limits::max(); // fp32 max is fp16 inf + migraphx::bf16 bf16_inf(fmax); + CHECK(bit_equal(bf16_inf, std::numeric_limits::max())); +} + +TEST_CASE(test_numeric_lowest_1) +{ + float flowest = std::numeric_limits::lowest(); + migraphx::bf16 bf16_neginf(flowest); + CHECK(bit_equal(bf16_neginf, std::numeric_limits::lowest())); +} + +TEST_CASE(test_max_eq_lowest) +{ + EXPECT(migraphx::float_equal(std::numeric_limits::lowest(), + -1 * std::numeric_limits::max())); +} + +TEST_CASE(test_isfinite) +{ + EXPECT(std::isfinite(migraphx::bf16(0.0))); + EXPECT(std::isfinite(migraphx::bf16(-0.0))); + EXPECT(not std::isfinite(migraphx::bf16(std::numeric_limits::quiet_NaN()))); +} + +TEST_CASE(test_binary_ops) +{ + auto a = migraphx::bf16(-1.0); + auto b = migraphx::bf16(1.0); + auto c = migraphx::bf16(0.0); + auto d = migraphx::bf16(-0.0); + EXPECT(migraphx::float_equal((c + d), c)); + EXPECT(migraphx::float_equal((c + d), d)); + EXPECT(migraphx::float_equal((a + b), c)); + EXPECT(migraphx::float_equal((a + b), d)); + + auto e = migraphx::bf16(10.0); + auto f = migraphx::bf16(-10.0); + EXPECT(e > f); + EXPECT(f < e); + EXPECT(f <= e); + EXPECT(e >= f); + EXPECT(e <= e); + EXPECT(f >= f); + EXPECT(not migraphx::float_equal(f, e)); +} + +TEST_CASE(test_stream_op) +{ + auto a = migraphx::bf16(-1.0); + std::stringstream ss; + ss << a; + EXPECT(std::string("-1") == ss.str()); + ss = std::stringstream(); + auto b = std::numeric_limits::quiet_NaN(); + ss << b; + EXPECT(std::string("nan") == ss.str()); +} + +int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/ref/add.cpp b/test/ref/add.cpp index ec9b7a2c40e..525e74dfa11 100644 --- a/test/ref/add.cpp +++ b/test/ref/add.cpp @@ -137,6 +137,25 @@ TEST_CASE(fp16_test) EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); } +TEST_CASE(bf16_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {1}}; + migraphx::bf16 a{1.5}; + migraphx::bf16 b{2.5}; + migraphx::bf16 c{4.0}; + auto l0 = mm->add_literal(migraphx::literal{s, {a}}); + auto l1 = mm->add_literal(migraphx::literal{s, {b}}); + mm->add_instruction(migraphx::make_op("add"), l0, l1); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + std::vector results_vector(1); + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{c}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + TEST_CASE(fp32_fp16_test) { auto create_program = [] { diff --git a/test/ref/isinf.cpp b/test/ref/isinf.cpp index 900ebc13fda..4d86e7f1b27 100644 --- a/test/ref/isinf.cpp +++ b/test/ref/isinf.cpp @@ -83,6 +83,25 @@ TEST_CASE(isinf_half_test) EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); } +TEST_CASE(isinf_bf16_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {2, 3}}; + auto inf_val = std::numeric_limits::infinity(); + migraphx::bf16 a{1.2}; + migraphx::bf16 b{5.2}; + std::vector data0 = {a, b, inf_val, -inf_val, b, a}; + auto l1 = mm->add_literal(migraphx::literal{s, data0}); + mm->add_instruction(migraphx::make_op("isinf"), l1); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold = {0, 0, 1, 1, 0, 0}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + TEST_CASE(isinf_dyn_test) { migraphx::program p; diff --git a/tools/api/migraphx.h b/tools/api/migraphx.h index 4ce7ea1d07b..df89fee6c37 100644 --- a/tools/api/migraphx.h +++ b/tools/api/migraphx.h @@ -47,7 +47,8 @@ m(uint64_type, uint64_t) \ m(fp8e4m3fnuz_type, migraphx::fp8::fp8e4m3fnuz) \ m(fp8e4m3fn_type, migraphx::fp8::fp8e4m3fn) \ - m(fp8e5m2_type, migraphx::fp8::fp8e5m2) + m(fp8e5m2_type, migraphx::fp8::fp8e5m2) \ + m(bf16_type, bf16) // clang-format on #ifdef __cplusplus