pocketpy/3rd/numpy/include/xtl/xhalf_float_impl.hpp

// half - IEEE 754-based half-precision floating-point library.
//
// Copyright (c) 2012-2019 Christian Rau <rauy@users.sourceforge.net>
// Copyright (c) 2020 0xBYTESHIFT
//
// 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.

/// \file
/// Main header file for half-precision functionality.

#pragma once

#define HALF_TWOS_COMPLEMENT_INT 1

// any error throwing C++ exceptions?
#if defined(HALF_ERRHANDLING_THROW_INVALID) || defined(HALF_ERRHANDLING_THROW_DIVBYZERO) || defined(HALF_ERRHANDLING_THROW_OVERFLOW) || defined(HALF_ERRHANDLING_THROW_UNDERFLOW) || defined(HALF_ERRHANDLING_THROW_INEXACT)
#define HALF_ERRHANDLING_THROWS 1
#endif

// any error handling enabled?
#define HALF_ERRHANDLING	(HALF_ERRHANDLING_FLAGS||HALF_ERRHANDLING_ERRNO||HALF_ERRHANDLING_FENV||HALF_ERRHANDLING_THROWS)

#if HALF_ERRHANDLING
	#define HALF_UNUSED_NOERR(name) name
#else
	#define HALF_UNUSED_NOERR(name)
#endif

// support constexpr
#if HALF_ERRHANDLING
	#define constexpr_NOERR
#else
	#define constexpr_NOERR	constexpr
#endif

#include <utility>
#include <algorithm>
#include <istream>
#include <ostream>
#include <limits>
#include <stdexcept>
#include <climits>
#include <cmath>
#include <cstring>
#include <cstdlib>
#include <type_traits>
#include <cstdint>
#if HALF_ERRHANDLING_ERRNO
	#include <cerrno>
#endif
#include <cfenv>
#include <functional>

#ifndef HALF_ENABLE_F16C_INTRINSICS
	/// Enable F16C intruction set intrinsics.
	/// Defining this to 1 enables the use of [F16C compiler intrinsics](https://en.wikipedia.org/wiki/F16C) for converting between 
	/// half-precision and single-precision values which may result in improved performance. This will not perform additional checks 
	/// for support of the F16C instruction set, so an appropriate target platform is required when enabling this feature.
	///
	/// Unless predefined it will be enabled automatically when the `__F16C__` symbol is defined, which some compilers do on supporting platforms.
	#define HALF_ENABLE_F16C_INTRINSICS __F16C__
#endif

#if HALF_ENABLE_F16C_INTRINSICS
	#include <immintrin.h>
#endif

#ifndef HALF_ERRHANDLING_OVERFLOW_TO_INEXACT
/// Raise INEXACT exception on overflow.
/// Defining this to 1 (default) causes overflow errors to automatically raise inexact exceptions in addition.
/// These will be raised after any possible handling of the underflow exception.
#define HALF_ERRHANDLING_OVERFLOW_TO_INEXACT	1
#endif

#ifndef HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT
/// Raise INEXACT exception on underflow.
/// Defining this to 1 (default) causes underflow errors to automatically raise inexact exceptions in addition.
/// These will be raised after any possible handling of the underflow exception.
///
/// **Note:** This will actually cause underflow (and the accompanying inexact) exceptions to be raised *only* when the result 
/// is inexact, while if disabled bare underflow errors will be raised for *any* (possibly exact) subnormal result.
#define HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT	1
#endif

/// Default rounding mode.
/// This specifies the rounding mode used for all conversions between [half](\ref half_float::half)s and more precise types 
/// (unless using half_cast() and specifying the rounding mode directly) as well as in arithmetic operations and mathematical 
/// functions. It can be redefined (before including half.hpp) to one of the standard rounding modes using their respective 
/// constants or the equivalent values of 
/// [std::float_round_style](https://en.cppreference.com/w/cpp/types/numeric_limits/float_round_style):
///
/// `std::float_round_style`         | value | rounding
/// ---------------------------------|-------|-------------------------
/// `std::round_indeterminate`       | -1    | fastest
/// `std::round_toward_zero`         | 0     | toward zero
/// `std::round_to_nearest`          | 1     | to nearest (default)
/// `std::round_toward_infinity`     | 2     | toward positive infinity
/// `std::round_toward_neg_infinity` | 3     | toward negative infinity
///
/// By default this is set to `1` (`std::round_to_nearest`), which rounds results to the nearest representable value. It can even 
/// be set to [std::numeric_limits<float>::round_style](https://en.cppreference.com/w/cpp/types/numeric_limits/round_style) to synchronize 
/// the rounding mode with that of the built-in single-precision implementation (which is likely `std::round_to_nearest`, though).
#ifndef HALF_ROUND_STYLE
	#define HALF_ROUND_STYLE	1		// = std::round_to_nearest
#endif

/// Value signaling overflow.
/// In correspondence with `HUGE_VAL[F|L]` from `<cmath>` this symbol expands to a positive value signaling the overflow of an 
/// operation, in particular it just evaluates to positive infinity.
///
/// **See also:** Documentation for [HUGE_VAL](https://en.cppreference.com/w/cpp/numeric/math/HUGE_VAL)
#define HUGE_VALH	std::numeric_limits<half_float::half>::infinity()

/// Fast half-precision fma function.
/// This symbol is defined if the fma() function generally executes as fast as, or faster than, a separate 
/// half-precision multiplication followed by an addition, which is always the case.
///
/// **See also:** Documentation for [FP_FAST_FMA](https://en.cppreference.com/w/cpp/numeric/math/fma)
#define FP_FAST_FMAH	1

///	Half rounding mode.
/// In correspondence with `FLT_ROUNDS` from `<cfloat>` this symbol expands to the rounding mode used for 
/// half-precision operations. It is an alias for [HALF_ROUND_STYLE](\ref HALF_ROUND_STYLE).
///
/// **See also:** Documentation for [FLT_ROUNDS](https://en.cppreference.com/w/cpp/types/climits/FLT_ROUNDS)
#define HLF_ROUNDS	HALF_ROUND_STYLE

#ifndef FP_ILOGB0
	#define FP_ILOGB0		INT_MIN
#endif
#ifndef FP_ILOGBNAN
	#define FP_ILOGBNAN		INT_MAX
#endif
#ifndef FP_SUBNORMAL
	#define FP_SUBNORMAL	0
#endif
#ifndef FP_ZERO
	#define FP_ZERO			1
#endif
#ifndef FP_NAN
	#define FP_NAN			2
#endif
#ifndef FP_INFINITE
	#define FP_INFINITE		3
#endif
#ifndef FP_NORMAL
	#define FP_NORMAL		4
#endif

#if !defined(FE_ALL_EXCEPT)
	#define FE_INVALID		0x10
	#define FE_DIVBYZERO	0x08
	#define FE_OVERFLOW		0x04
	#define FE_UNDERFLOW	0x02
	#define FE_INEXACT		0x01
	#define FE_ALL_EXCEPT	(FE_INVALID|FE_DIVBYZERO|FE_OVERFLOW|FE_UNDERFLOW|FE_INEXACT)
#endif


/// Main namespace for half-precision functionality.
/// This namespace contains all the functionality provided by the library.
namespace half_float {
	class half;

	/// Library-defined half-precision literals.
	/// Import this namespace to enable half-precision floating-point literals:
	/// ~~~~{.cpp}
	/// using namespace half_float::literal;
	/// half_float::half = 4.2_h;
	/// ~~~~
	namespace literal {
		half operator "" _h(long double);
	}

	/// \internal
	/// \brief Implementation details.
	namespace detail {
		/// Conditional type.
		template<bool B, class T, class F> struct conditional : std::conditional<B,T,F> {};

		/// Helper for tag dispatching.
		template<bool B> struct bool_type : std::integral_constant<bool,B> {};
		using std::true_type;
		using std::false_type;

		/// Type traits for floating-point types.
		template<class T> struct is_float : std::is_floating_point<T> {};

		/// Type traits for floating-point bits.
		template<class T> struct bits { using type = unsigned char; };
		template<class T> struct bits<const T> : bits<T> {};
		template<class T> struct bits<volatile T> : bits<T> {};
		template<class T> struct bits<const volatile T> : bits<T> {};

		/// Unsigned integer of (at least) 16 bits width.
		using uint16 = std::uint_least16_t;

		/// Fastest unsigned integer of (at least) 32 bits width.
		using uint32 = std::uint_fast32_t;

		/// Fastest signed integer of (at least) 32 bits width.
		using int32 = std::int_fast32_t;

		/// Unsigned integer of (at least) 32 bits width.
		template<> struct bits<float> { using type = std::uint_least32_t; };

		/// Unsigned integer of (at least) 64 bits width.
		template<> struct bits<double> { using type = std::uint_least64_t; };
		template<class T> using bits_t = typename bits<T>::type;

	#ifdef HALF_ARITHMETIC_TYPE
		/// Type to use for arithmetic computations and mathematic functions internally.
		typedef HALF_ARITHMETIC_TYPE internal_t;
	#endif

		/// Tag type for binary construction.
		struct binary_t {};

		/// Tag for binary construction.
		constexpr binary_t binary = binary_t();

		/// \name Implementation defined classification and arithmetic
		/// \{

		/// Check for infinity.
		/// \tparam T argument type (builtin floating-point type)
		/// \param arg value to query
		/// \retval true if infinity
		/// \retval false else
		template<class T> bool builtin_isinf(T arg) { return std::isinf(arg); }

		/// Check for NaN.
		/// \tparam T argument type (builtin floating-point type)
		/// \param arg value to query
		/// \retval true if not a number
		/// \retval false else
		template<class T> bool builtin_isnan(T arg) { return std::isnan(arg); }

		/// Check sign.
		/// \tparam T argument type (builtin floating-point type)
		/// \param arg value to query
		/// \retval true if signbit set
		/// \retval false else
		template<class T> bool builtin_signbit(T arg) { return std::signbit(arg); }

		/// Platform-independent sign mask.
		/// \param arg integer value in two's complement
		/// \retval -1 if \a arg negative
		/// \retval 0 if \a arg positive
		inline uint32 sign_mask(uint32 arg) {
			static const int N = std::numeric_limits<uint32>::digits - 1;
		#if HALF_TWOS_COMPLEMENT_INT
			return static_cast<int32>(arg) >> N;
		#else
			return -((arg>>N)&1);
		#endif
		}

		/// Platform-independent arithmetic right shift.
		/// \param arg integer value in two's complement
		/// \param i shift amount (at most 31)
		/// \return \a arg right shifted for \a i bits with possible sign extension
		inline uint32 arithmetic_shift(uint32 arg, int i) {
		#if HALF_TWOS_COMPLEMENT_INT
			return static_cast<int32>(arg) >> i;
		#else
			return static_cast<int32>(arg)/(static_cast<int32>(1)<<i) - ((arg>>(std::numeric_limits<uint32>::digits-1))&1);
		#endif
		}

		/// \}
		/// \name Error handling
		/// \{

		/// Internal exception flags.
		/// \return reference to global exception flags
		inline int& errflags() { thread_local int flags = 0; return flags; }

		/// Raise floating-point exception.
		/// \param flags exceptions to raise
		/// \param cond condition to raise exceptions for
		inline void raise(int HALF_UNUSED_NOERR(flags), bool HALF_UNUSED_NOERR(cond) = true) {
		#if HALF_ERRHANDLING
			if(!cond)
				return;
		#if HALF_ERRHANDLING_FLAGS
			errflags() |= flags;
		#endif
		#if HALF_ERRHANDLING_ERRNO
			if(flags & FE_INVALID)
				errno = EDOM;
			else if(flags & (FE_DIVBYZERO|FE_OVERFLOW|FE_UNDERFLOW))
				errno = ERANGE;
		#endif
		#if HALF_ERRHANDLING_FENV
			std::feraiseexcept(flags);
		#endif
		#ifdef HALF_ERRHANDLING_THROW_INVALID
			if(flags & FE_INVALID)
				throw std::domain_error(HALF_ERRHANDLING_THROW_INVALID);
		#endif
		#ifdef HALF_ERRHANDLING_THROW_DIVBYZERO
			if(flags & FE_DIVBYZERO)
				throw std::domain_error(HALF_ERRHANDLING_THROW_DIVBYZERO);
		#endif
		#ifdef HALF_ERRHANDLING_THROW_OVERFLOW
			if(flags & FE_OVERFLOW)
				throw std::overflow_error(HALF_ERRHANDLING_THROW_OVERFLOW);
		#endif
		#ifdef HALF_ERRHANDLING_THROW_UNDERFLOW
			if(flags & FE_UNDERFLOW)
				throw std::underflow_error(HALF_ERRHANDLING_THROW_UNDERFLOW);
		#endif
		#ifdef HALF_ERRHANDLING_THROW_INEXACT
			if(flags & FE_INEXACT)
				throw std::range_error(HALF_ERRHANDLING_THROW_INEXACT);
		#endif
		#if HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT
			if((flags & FE_UNDERFLOW) && !(flags & FE_INEXACT))
				raise(FE_INEXACT);
		#endif
		#if HALF_ERRHANDLING_OVERFLOW_TO_INEXACT
			if((flags & FE_OVERFLOW) && !(flags & FE_INEXACT))
				raise(FE_INEXACT);
		#endif
		#endif
		}

		/// Check and signal for any NaN.
		/// \param x first half-precision value to check
		/// \param y second half-precision value to check
		/// \retval true if either \a x or \a y is NaN
		/// \retval false else
		/// \exception FE_INVALID if \a x or \a y is NaN
		inline constexpr_NOERR bool compsignal(unsigned int x, unsigned int y) {
		#if HALF_ERRHANDLING
			raise(FE_INVALID, (x&0x7FFF)>0x7C00 || (y&0x7FFF)>0x7C00);
		#endif
			return (x&0x7FFF) > 0x7C00 || (y&0x7FFF) > 0x7C00;
		}

		/// Signal and silence signaling NaN.
		/// \param nan half-precision NaN value
		/// \return quiet NaN
		/// \exception FE_INVALID if \a nan is signaling NaN
		inline constexpr_NOERR unsigned int signal(unsigned int nan) {
		#if HALF_ERRHANDLING
			raise(FE_INVALID, !(nan&0x200));
		#endif
			return nan | 0x200;
		}

		/// Signal and silence signaling NaNs.
		/// \param x first half-precision value to check
		/// \param y second half-precision value to check
		/// \return quiet NaN
		/// \exception FE_INVALID if \a x or \a y is signaling NaN
		inline constexpr_NOERR unsigned int signal(unsigned int x, unsigned int y) {
		#if HALF_ERRHANDLING
			raise(FE_INVALID, ((x&0x7FFF)>0x7C00 && !(x&0x200)) || ((y&0x7FFF)>0x7C00 && !(y&0x200)));
		#endif
			return ((x&0x7FFF)>0x7C00) ? (x|0x200) : (y|0x200);
		}

		/// Signal and silence signaling NaNs.
		/// \param x first half-precision value to check
		/// \param y second half-precision value to check
		/// \param z third half-precision value to check
		/// \return quiet NaN
		/// \exception FE_INVALID if \a x, \a y or \a z is signaling NaN
		inline constexpr_NOERR unsigned int signal(unsigned int x, unsigned int y, unsigned int z) {
		#if HALF_ERRHANDLING
			raise(FE_INVALID, ((x&0x7FFF)>0x7C00 && !(x&0x200)) || ((y&0x7FFF)>0x7C00 && !(y&0x200)) || ((z&0x7FFF)>0x7C00 && !(z&0x200)));
		#endif
			return ((x&0x7FFF)>0x7C00) ? (x|0x200) : ((y&0x7FFF)>0x7C00) ? (y|0x200) : (z|0x200);
		}

		/// Select value or signaling NaN.
		/// \param x preferred half-precision value
		/// \param y ignored half-precision value except for signaling NaN
		/// \return \a y if signaling NaN, \a x otherwise
		/// \exception FE_INVALID if \a y is signaling NaN
		inline constexpr_NOERR unsigned int select(unsigned int x, unsigned int HALF_UNUSED_NOERR(y)) {
		#if HALF_ERRHANDLING
			return (((y&0x7FFF)>0x7C00) && !(y&0x200)) ? signal(y) : x;
		#else
			return x;
		#endif
		}

		/// Raise domain error and return NaN.
		/// return quiet NaN
		/// \exception FE_INVALID
		inline constexpr_NOERR unsigned int invalid() {
		#if HALF_ERRHANDLING
			raise(FE_INVALID);
		#endif
			return 0x7FFF;
		}

		/// Raise pole error and return infinity.
		/// \param sign half-precision value with sign bit only
		/// \return half-precision infinity with sign of \a sign
		/// \exception FE_DIVBYZERO
		inline constexpr_NOERR unsigned int pole(unsigned int sign = 0) {
		#if HALF_ERRHANDLING
			raise(FE_DIVBYZERO);
		#endif
			return sign | 0x7C00;
		}

		/// Check value for underflow.
		/// \param arg non-zero half-precision value to check
		/// \return \a arg
		/// \exception FE_UNDERFLOW if arg is subnormal
		inline constexpr_NOERR unsigned int check_underflow(unsigned int arg) {
		#if HALF_ERRHANDLING && !HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT
			raise(FE_UNDERFLOW, !(arg&0x7C00));
		#endif
			return arg;
		}

		/// \}
		/// \name Conversion and rounding
		/// \{

		/// Half-precision overflow.
		/// \tparam R rounding mode to use
		/// \param sign half-precision value with sign bit only
		/// \return rounded overflowing half-precision value
		/// \exception FE_OVERFLOW
		template<std::float_round_style R> constexpr_NOERR unsigned int overflow(unsigned int sign = 0) {
		#if HALF_ERRHANDLING
			raise(FE_OVERFLOW);
		#endif
			return	(R==std::round_toward_infinity) ? (sign+0x7C00-(sign>>15)) :
					(R==std::round_toward_neg_infinity) ? (sign+0x7BFF+(sign>>15)) :
					(R==std::round_toward_zero) ? (sign|0x7BFF) :
					(sign|0x7C00);
		}

		/// Half-precision underflow.
		/// \tparam R rounding mode to use
		/// \param sign half-precision value with sign bit only
		/// \return rounded underflowing half-precision value
		/// \exception FE_UNDERFLOW
		template<std::float_round_style R> constexpr_NOERR unsigned int underflow(unsigned int sign = 0) {
		#if HALF_ERRHANDLING
			raise(FE_UNDERFLOW);
		#endif
			return	(R==std::round_toward_infinity) ? (sign+1-(sign>>15)) :
					(R==std::round_toward_neg_infinity) ? (sign+(sign>>15)) :
					sign;
		}

		/// Round half-precision number.
		/// \tparam R rounding mode to use
		/// \tparam I `true` to always raise INEXACT exception, `false` to raise only for rounded results
		/// \param value finite half-precision number to round
		/// \param g guard bit (most significant discarded bit)
		/// \param s sticky bit (or of all but the most significant discarded bits)
		/// \return rounded half-precision value
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded or \a I is `true`
		template<std::float_round_style R,bool I> constexpr_NOERR unsigned int rounded(unsigned int value, int g, int s) {
		#if HALF_ERRHANDLING
			value +=	(R==std::round_to_nearest) ? (g&(s|value)) :
						(R==std::round_toward_infinity) ? (~(value>>15)&(g|s)) :
						(R==std::round_toward_neg_infinity) ? ((value>>15)&(g|s)) : 0;
			if((value&0x7C00) == 0x7C00)
				raise(FE_OVERFLOW);
			else if(value & 0x7C00)
				raise(FE_INEXACT, I || (g|s)!=0);
			else
				raise(FE_UNDERFLOW, !(HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT) || I || (g|s)!=0);
			return value;
		#else
			return	(R==std::round_to_nearest) ? (value+(g&(s|value))) :
					(R==std::round_toward_infinity) ? (value+(~(value>>15)&(g|s))) :
					(R==std::round_toward_neg_infinity) ? (value+((value>>15)&(g|s))) :
					value;
		#endif
		}

		/// Round half-precision number to nearest integer value.
		/// \tparam R rounding mode to use
		/// \tparam E `true` for round to even, `false` for round away from zero
		/// \tparam I `true` to raise INEXACT exception (if inexact), `false` to never raise it
		/// \param value half-precision value to round
		/// \return half-precision bits for nearest integral value
		/// \exception FE_INVALID for signaling NaN
		/// \exception FE_INEXACT if value had to be rounded and \a I is `true`
		template<std::float_round_style R,bool E,bool I> unsigned int integral(unsigned int value) {
			unsigned int abs = value & 0x7FFF;
			if(abs < 0x3C00) {
				raise(FE_INEXACT, I);
				return ((R==std::round_to_nearest) ? (0x3C00&-static_cast<unsigned>(abs>=(0x3800+E))) :
						(R==std::round_toward_infinity) ? (0x3C00&-(~(value>>15)&(abs!=0))) :
						(R==std::round_toward_neg_infinity) ? (0x3C00&-static_cast<unsigned>(value>0x8000)) :
						0) | (value&0x8000);
			}
			if(abs >= 0x6400)
				return (abs>0x7C00) ? signal(value) : value;
			unsigned int exp = 25 - (abs>>10), mask = (1<<exp) - 1;
			raise(FE_INEXACT, I && (value&mask));
			return ((	(R==std::round_to_nearest) ? ((1<<(exp-1))-(~(value>>exp)&E)) :
						(R==std::round_toward_infinity) ? (mask&((value>>15)-1)) :
						(R==std::round_toward_neg_infinity) ? (mask&-(value>>15)) :
						0) + value) & ~mask;
		}

		/// Convert fixed point to half-precision floating-point.
		/// \tparam R rounding mode to use
		/// \tparam F number of fractional bits (at least 11)
		/// \tparam S `true` for signed, `false` for unsigned
		/// \tparam N `true` for additional normalization step, `false` if already normalized to 1.F
		/// \tparam I `true` to always raise INEXACT exception, `false` to raise only for rounded results
		/// \param m mantissa in Q1.F fixed point format
		/// \param exp exponent
		/// \param sign half-precision value with sign bit only
		/// \param s sticky bit (or of all but the most significant already discarded bits)
		/// \return value converted to half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded or \a I is `true`
		template<std::float_round_style R,unsigned int F,bool S,bool N,bool I> unsigned int fixed2half(uint32 m, int exp = 14, unsigned int sign = 0, int s = 0) {
			if(S) {
				uint32 msign = sign_mask(m);
				m = (m^msign) - msign;
				sign = msign & 0x8000;
			}
			if(N)
				for(; m<(static_cast<uint32>(1)<<F) && exp; m<<=1,--exp) ;
			else if(exp < 0)
				return rounded<R,I>(sign+(m>>(F-10-exp)), (m>>(F-11-exp))&1, s|((m&((static_cast<uint32>(1)<<(F-11-exp))-1))!=0));
			return rounded<R,I>(sign+(exp<<10)+(m>>(F-10)), (m>>(F-11))&1, s|((m&((static_cast<uint32>(1)<<(F-11))-1))!=0));
		}

		/// Convert IEEE single-precision to half-precision.
		/// Credit for this goes to [Jeroen van der Zijp](ftp://ftp.fox-toolkit.org/pub/fasthalffloatconversion.pdf).
		/// \tparam R rounding mode to use
		/// \param value single-precision value to convert
		/// \return rounded half-precision value
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded
		template<std::float_round_style R> unsigned int float2half_impl(float value, true_type) {
		#if HALF_ENABLE_F16C_INTRINSICS
			return _mm_cvtsi128_si32(_mm_cvtps_ph(_mm_set_ss(value),
				(R==std::round_to_nearest) ? _MM_FROUND_TO_NEAREST_INT :
				(R==std::round_toward_zero) ? _MM_FROUND_TO_ZERO :
				(R==std::round_toward_infinity) ? _MM_FROUND_TO_POS_INF :
				(R==std::round_toward_neg_infinity) ? _MM_FROUND_TO_NEG_INF :
				_MM_FROUND_CUR_DIRECTION));
		#else
			bits_t<float> fbits;
			std::memcpy(&fbits, &value, sizeof(float));
		#if 1
			unsigned int sign = (fbits>>16) & 0x8000;
			fbits &= 0x7FFFFFFF;
			if(fbits >= 0x7F800000)
				return sign | 0x7C00 | ((fbits>0x7F800000) ? (0x200|((fbits>>13)&0x3FF)) : 0);
			if(fbits >= 0x47800000)
				return overflow<R>(sign);
			if(fbits >= 0x38800000)
				return rounded<R,false>(sign|(((fbits>>23)-112)<<10)|((fbits>>13)&0x3FF), (fbits>>12)&1, (fbits&0xFFF)!=0);
			if(fbits >= 0x33000000)
			{
				int i = 125 - (fbits>>23);
				fbits = (fbits&0x7FFFFF) | 0x800000;
				return rounded<R,false>(sign|(fbits>>(i+1)), (fbits>>i)&1, (fbits&((static_cast<uint32>(1)<<i)-1))!=0);
			}
			if(fbits != 0)
				return underflow<R>(sign);
			return sign;
		#else
			static const uint16 base_table[512] = {
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 
				0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0000, 0x0001, 0x0002, 0x0004, 0x0008, 0x0010, 0x0020, 0x0040, 0x0080, 0x0100, 
				0x0200, 0x0400, 0x0800, 0x0C00, 0x1000, 0x1400, 0x1800, 0x1C00, 0x2000, 0x2400, 0x2800, 0x2C00, 0x3000, 0x3400, 0x3800, 0x3C00, 
				0x4000, 0x4400, 0x4800, 0x4C00, 0x5000, 0x5400, 0x5800, 0x5C00, 0x6000, 0x6400, 0x6800, 0x6C00, 0x7000, 0x7400, 0x7800, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 
				0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7BFF, 0x7C00, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 
				0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8000, 0x8001, 0x8002, 0x8004, 0x8008, 0x8010, 0x8020, 0x8040, 0x8080, 0x8100, 
				0x8200, 0x8400, 0x8800, 0x8C00, 0x9000, 0x9400, 0x9800, 0x9C00, 0xA000, 0xA400, 0xA800, 0xAC00, 0xB000, 0xB400, 0xB800, 0xBC00, 
				0xC000, 0xC400, 0xC800, 0xCC00, 0xD000, 0xD400, 0xD800, 0xDC00, 0xE000, 0xE400, 0xE800, 0xEC00, 0xF000, 0xF400, 0xF800, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 
				0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFBFF, 0xFC00 };
			static const unsigned char shift_table[256] = {
				24, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 
				25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 
				25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 25, 
				25, 25, 25, 25, 25, 25, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 
				13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 13, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 
				24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 
				24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 
				24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 13 };
			int sexp = fbits >> 23, exp = sexp & 0xFF, i = shift_table[exp];
			fbits &= 0x7FFFFF;
			uint32 m = (fbits|((exp!=0)<<23)) & -static_cast<uint32>(exp!=0xFF);
			return rounded<R,false>(base_table[sexp]+(fbits>>i), (m>>(i-1))&1, (((static_cast<uint32>(1)<<(i-1))-1)&m)!=0);
		#endif
		#endif
		}

		/// Convert IEEE double-precision to half-precision.
		/// \tparam R rounding mode to use
		/// \param value double-precision value to convert
		/// \return rounded half-precision value
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded
		template<std::float_round_style R> unsigned int float2half_impl(double value, true_type) {
		#if HALF_ENABLE_F16C_INTRINSICS
			if(R == std::round_indeterminate)
				return _mm_cvtsi128_si32(_mm_cvtps_ph(_mm_cvtpd_ps(_mm_set_sd(value)), _MM_FROUND_CUR_DIRECTION));
		#endif
			bits_t<double> dbits;
			std::memcpy(&dbits, &value, sizeof(double));
			uint32 hi = dbits >> 32, lo = dbits & 0xFFFFFFFF;
			unsigned int sign = (hi>>16) & 0x8000;
			hi &= 0x7FFFFFFF;
			if(hi >= 0x7FF00000)
				return sign | 0x7C00 | ((dbits&0xFFFFFFFFFFFFF) ? (0x200|((hi>>10)&0x3FF)) : 0);
			if(hi >= 0x40F00000)
				return overflow<R>(sign);
			if(hi >= 0x3F100000)
				return rounded<R,false>(sign|(((hi>>20)-1008)<<10)|((hi>>10)&0x3FF), (hi>>9)&1, ((hi&0x1FF)|lo)!=0);
			if(hi >= 0x3E600000) {
				int i = 1018 - (hi>>20);
				hi = (hi&0xFFFFF) | 0x100000;
				return rounded<R,false>(sign|(hi>>(i+1)), (hi>>i)&1, ((hi&((static_cast<uint32>(1)<<i)-1))|lo)!=0);
			}
			if((hi|lo) != 0)
				return underflow<R>(sign);
			return sign;
		}

		/// Convert non-IEEE floating-point to half-precision.
		/// \tparam R rounding mode to use
		/// \tparam T source type (builtin floating-point type)
		/// \param value floating-point value to convert
		/// \return rounded half-precision value
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded
		template<std::float_round_style R,class T> unsigned int float2half_impl(T value, ...) {
			unsigned int hbits = static_cast<unsigned>(builtin_signbit(value)) << 15;
			if(value == T())
				return hbits;
			if(builtin_isnan(value))
				return hbits | 0x7FFF;
			if(builtin_isinf(value))
				return hbits | 0x7C00;
			int exp;
			std::frexp(value, &exp);
			if(exp > 16)
				return overflow<R>(hbits);
			if(exp < -13)
				value = std::ldexp(value, 25);
			else {
				value = std::ldexp(value, 12-exp);
				hbits |= ((exp+13)<<10);
			}
			T ival, frac = std::modf(value, &ival);
			int m = std::abs(static_cast<int>(ival));
			return rounded<R,false>(hbits+(m>>1), m&1, frac!=T());
		}

		/// Convert floating-point to half-precision.
		/// \tparam R rounding mode to use
		/// \tparam T source type (builtin floating-point type)
		/// \param value floating-point value to convert
		/// \return rounded half-precision value
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded
		template<std::float_round_style R,class T> unsigned int float2half(T value) {
			return float2half_impl<R>(value, bool_type<std::numeric_limits<T>::is_iec559&&sizeof(bits_t<T>)==sizeof(T)>());
		}
		template<class T> unsigned int float2half(T value) {
			return float2half_impl<(std::float_round_style)(HALF_ROUND_STYLE)>(value, bool_type<std::numeric_limits<T>::is_iec559&&sizeof(bits_t<T>)==sizeof(T)>());
		}

		/// Convert integer to half-precision floating-point.
		/// \tparam R rounding mode to use
		/// \tparam T type to convert (builtin integer type)
		/// \param value integral value to convert
		/// \return rounded half-precision value
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_INEXACT if value had to be rounded
		template<std::float_round_style R,class T> unsigned int int2half(T value) {
			unsigned int bits = static_cast<unsigned>(value<0) << 15;
			if(!value)
				return bits;
			if(bits)
				value = -value;
			if(value > 0xFFFF)
				return overflow<R>(bits);
			unsigned int m = static_cast<unsigned int>(value), exp = 24;
			for(; m<0x400; m<<=1,--exp) ;
			for(; m>0x7FF; m>>=1,++exp) ;
			bits |= (exp<<10) + m;
			return (exp>24) ? rounded<R,false>(bits, (value>>(exp-25))&1, (((1<<(exp-25))-1)&value)!=0) : bits;
		}

		/// Convert half-precision to IEEE single-precision.
		/// Credit for this goes to [Jeroen van der Zijp](ftp://ftp.fox-toolkit.org/pub/fasthalffloatconversion.pdf).
		/// \param value half-precision value to convert
		/// \return single-precision value
		inline float half2float_impl(unsigned int value, float, true_type) {
		#if HALF_ENABLE_F16C_INTRINSICS
			return _mm_cvtss_f32(_mm_cvtph_ps(_mm_cvtsi32_si128(value)));
		#else
		#if 0
			bits_t<float> fbits = static_cast<bits_t<float>>(value&0x8000) << 16;
			int abs = value & 0x7FFF;
			if(abs)
			{
				fbits |= 0x38000000 << static_cast<unsigned>(abs>=0x7C00);
				for(; abs<0x400; abs<<=1,fbits-=0x800000) ;
				fbits += static_cast<bits_t<float>>(abs) << 13;
			}
		#else
			static const bits_t<float> mantissa_table[2048] = {
				0x00000000, 0x33800000, 0x34000000, 0x34400000, 0x34800000, 0x34A00000, 0x34C00000, 0x34E00000, 0x35000000, 0x35100000, 0x35200000, 0x35300000, 0x35400000, 0x35500000, 0x35600000, 0x35700000, 
				0x35800000, 0x35880000, 0x35900000, 0x35980000, 0x35A00000, 0x35A80000, 0x35B00000, 0x35B80000, 0x35C00000, 0x35C80000, 0x35D00000, 0x35D80000, 0x35E00000, 0x35E80000, 0x35F00000, 0x35F80000, 
				0x36000000, 0x36040000, 0x36080000, 0x360C0000, 0x36100000, 0x36140000, 0x36180000, 0x361C0000, 0x36200000, 0x36240000, 0x36280000, 0x362C0000, 0x36300000, 0x36340000, 0x36380000, 0x363C0000, 
				0x36400000, 0x36440000, 0x36480000, 0x364C0000, 0x36500000, 0x36540000, 0x36580000, 0x365C0000, 0x36600000, 0x36640000, 0x36680000, 0x366C0000, 0x36700000, 0x36740000, 0x36780000, 0x367C0000, 
				0x36800000, 0x36820000, 0x36840000, 0x36860000, 0x36880000, 0x368A0000, 0x368C0000, 0x368E0000, 0x36900000, 0x36920000, 0x36940000, 0x36960000, 0x36980000, 0x369A0000, 0x369C0000, 0x369E0000, 
				0x36A00000, 0x36A20000, 0x36A40000, 0x36A60000, 0x36A80000, 0x36AA0000, 0x36AC0000, 0x36AE0000, 0x36B00000, 0x36B20000, 0x36B40000, 0x36B60000, 0x36B80000, 0x36BA0000, 0x36BC0000, 0x36BE0000, 
				0x36C00000, 0x36C20000, 0x36C40000, 0x36C60000, 0x36C80000, 0x36CA0000, 0x36CC0000, 0x36CE0000, 0x36D00000, 0x36D20000, 0x36D40000, 0x36D60000, 0x36D80000, 0x36DA0000, 0x36DC0000, 0x36DE0000, 
				0x36E00000, 0x36E20000, 0x36E40000, 0x36E60000, 0x36E80000, 0x36EA0000, 0x36EC0000, 0x36EE0000, 0x36F00000, 0x36F20000, 0x36F40000, 0x36F60000, 0x36F80000, 0x36FA0000, 0x36FC0000, 0x36FE0000, 
				0x37000000, 0x37010000, 0x37020000, 0x37030000, 0x37040000, 0x37050000, 0x37060000, 0x37070000, 0x37080000, 0x37090000, 0x370A0000, 0x370B0000, 0x370C0000, 0x370D0000, 0x370E0000, 0x370F0000, 
				0x37100000, 0x37110000, 0x37120000, 0x37130000, 0x37140000, 0x37150000, 0x37160000, 0x37170000, 0x37180000, 0x37190000, 0x371A0000, 0x371B0000, 0x371C0000, 0x371D0000, 0x371E0000, 0x371F0000, 
				0x37200000, 0x37210000, 0x37220000, 0x37230000, 0x37240000, 0x37250000, 0x37260000, 0x37270000, 0x37280000, 0x37290000, 0x372A0000, 0x372B0000, 0x372C0000, 0x372D0000, 0x372E0000, 0x372F0000, 
				0x37300000, 0x37310000, 0x37320000, 0x37330000, 0x37340000, 0x37350000, 0x37360000, 0x37370000, 0x37380000, 0x37390000, 0x373A0000, 0x373B0000, 0x373C0000, 0x373D0000, 0x373E0000, 0x373F0000, 
				0x37400000, 0x37410000, 0x37420000, 0x37430000, 0x37440000, 0x37450000, 0x37460000, 0x37470000, 0x37480000, 0x37490000, 0x374A0000, 0x374B0000, 0x374C0000, 0x374D0000, 0x374E0000, 0x374F0000, 
				0x37500000, 0x37510000, 0x37520000, 0x37530000, 0x37540000, 0x37550000, 0x37560000, 0x37570000, 0x37580000, 0x37590000, 0x375A0000, 0x375B0000, 0x375C0000, 0x375D0000, 0x375E0000, 0x375F0000, 
				0x37600000, 0x37610000, 0x37620000, 0x37630000, 0x37640000, 0x37650000, 0x37660000, 0x37670000, 0x37680000, 0x37690000, 0x376A0000, 0x376B0000, 0x376C0000, 0x376D0000, 0x376E0000, 0x376F0000, 
				0x37700000, 0x37710000, 0x37720000, 0x37730000, 0x37740000, 0x37750000, 0x37760000, 0x37770000, 0x37780000, 0x37790000, 0x377A0000, 0x377B0000, 0x377C0000, 0x377D0000, 0x377E0000, 0x377F0000, 
				0x37800000, 0x37808000, 0x37810000, 0x37818000, 0x37820000, 0x37828000, 0x37830000, 0x37838000, 0x37840000, 0x37848000, 0x37850000, 0x37858000, 0x37860000, 0x37868000, 0x37870000, 0x37878000, 
				0x37880000, 0x37888000, 0x37890000, 0x37898000, 0x378A0000, 0x378A8000, 0x378B0000, 0x378B8000, 0x378C0000, 0x378C8000, 0x378D0000, 0x378D8000, 0x378E0000, 0x378E8000, 0x378F0000, 0x378F8000, 
				0x37900000, 0x37908000, 0x37910000, 0x37918000, 0x37920000, 0x37928000, 0x37930000, 0x37938000, 0x37940000, 0x37948000, 0x37950000, 0x37958000, 0x37960000, 0x37968000, 0x37970000, 0x37978000, 
				0x37980000, 0x37988000, 0x37990000, 0x37998000, 0x379A0000, 0x379A8000, 0x379B0000, 0x379B8000, 0x379C0000, 0x379C8000, 0x379D0000, 0x379D8000, 0x379E0000, 0x379E8000, 0x379F0000, 0x379F8000, 
				0x37A00000, 0x37A08000, 0x37A10000, 0x37A18000, 0x37A20000, 0x37A28000, 0x37A30000, 0x37A38000, 0x37A40000, 0x37A48000, 0x37A50000, 0x37A58000, 0x37A60000, 0x37A68000, 0x37A70000, 0x37A78000, 
				0x37A80000, 0x37A88000, 0x37A90000, 0x37A98000, 0x37AA0000, 0x37AA8000, 0x37AB0000, 0x37AB8000, 0x37AC0000, 0x37AC8000, 0x37AD0000, 0x37AD8000, 0x37AE0000, 0x37AE8000, 0x37AF0000, 0x37AF8000, 
				0x37B00000, 0x37B08000, 0x37B10000, 0x37B18000, 0x37B20000, 0x37B28000, 0x37B30000, 0x37B38000, 0x37B40000, 0x37B48000, 0x37B50000, 0x37B58000, 0x37B60000, 0x37B68000, 0x37B70000, 0x37B78000, 
				0x37B80000, 0x37B88000, 0x37B90000, 0x37B98000, 0x37BA0000, 0x37BA8000, 0x37BB0000, 0x37BB8000, 0x37BC0000, 0x37BC8000, 0x37BD0000, 0x37BD8000, 0x37BE0000, 0x37BE8000, 0x37BF0000, 0x37BF8000, 
				0x37C00000, 0x37C08000, 0x37C10000, 0x37C18000, 0x37C20000, 0x37C28000, 0x37C30000, 0x37C38000, 0x37C40000, 0x37C48000, 0x37C50000, 0x37C58000, 0x37C60000, 0x37C68000, 0x37C70000, 0x37C78000, 
				0x37C80000, 0x37C88000, 0x37C90000, 0x37C98000, 0x37CA0000, 0x37CA8000, 0x37CB0000, 0x37CB8000, 0x37CC0000, 0x37CC8000, 0x37CD0000, 0x37CD8000, 0x37CE0000, 0x37CE8000, 0x37CF0000, 0x37CF8000, 
				0x37D00000, 0x37D08000, 0x37D10000, 0x37D18000, 0x37D20000, 0x37D28000, 0x37D30000, 0x37D38000, 0x37D40000, 0x37D48000, 0x37D50000, 0x37D58000, 0x37D60000, 0x37D68000, 0x37D70000, 0x37D78000, 
				0x37D80000, 0x37D88000, 0x37D90000, 0x37D98000, 0x37DA0000, 0x37DA8000, 0x37DB0000, 0x37DB8000, 0x37DC0000, 0x37DC8000, 0x37DD0000, 0x37DD8000, 0x37DE0000, 0x37DE8000, 0x37DF0000, 0x37DF8000, 
				0x37E00000, 0x37E08000, 0x37E10000, 0x37E18000, 0x37E20000, 0x37E28000, 0x37E30000, 0x37E38000, 0x37E40000, 0x37E48000, 0x37E50000, 0x37E58000, 0x37E60000, 0x37E68000, 0x37E70000, 0x37E78000, 
				0x37E80000, 0x37E88000, 0x37E90000, 0x37E98000, 0x37EA0000, 0x37EA8000, 0x37EB0000, 0x37EB8000, 0x37EC0000, 0x37EC8000, 0x37ED0000, 0x37ED8000, 0x37EE0000, 0x37EE8000, 0x37EF0000, 0x37EF8000, 
				0x37F00000, 0x37F08000, 0x37F10000, 0x37F18000, 0x37F20000, 0x37F28000, 0x37F30000, 0x37F38000, 0x37F40000, 0x37F48000, 0x37F50000, 0x37F58000, 0x37F60000, 0x37F68000, 0x37F70000, 0x37F78000, 
				0x37F80000, 0x37F88000, 0x37F90000, 0x37F98000, 0x37FA0000, 0x37FA8000, 0x37FB0000, 0x37FB8000, 0x37FC0000, 0x37FC8000, 0x37FD0000, 0x37FD8000, 0x37FE0000, 0x37FE8000, 0x37FF0000, 0x37FF8000, 
				0x38000000, 0x38004000, 0x38008000, 0x3800C000, 0x38010000, 0x38014000, 0x38018000, 0x3801C000, 0x38020000, 0x38024000, 0x38028000, 0x3802C000, 0x38030000, 0x38034000, 0x38038000, 0x3803C000, 
				0x38040000, 0x38044000, 0x38048000, 0x3804C000, 0x38050000, 0x38054000, 0x38058000, 0x3805C000, 0x38060000, 0x38064000, 0x38068000, 0x3806C000, 0x38070000, 0x38074000, 0x38078000, 0x3807C000, 
				0x38080000, 0x38084000, 0x38088000, 0x3808C000, 0x38090000, 0x38094000, 0x38098000, 0x3809C000, 0x380A0000, 0x380A4000, 0x380A8000, 0x380AC000, 0x380B0000, 0x380B4000, 0x380B8000, 0x380BC000, 
				0x380C0000, 0x380C4000, 0x380C8000, 0x380CC000, 0x380D0000, 0x380D4000, 0x380D8000, 0x380DC000, 0x380E0000, 0x380E4000, 0x380E8000, 0x380EC000, 0x380F0000, 0x380F4000, 0x380F8000, 0x380FC000, 
				0x38100000, 0x38104000, 0x38108000, 0x3810C000, 0x38110000, 0x38114000, 0x38118000, 0x3811C000, 0x38120000, 0x38124000, 0x38128000, 0x3812C000, 0x38130000, 0x38134000, 0x38138000, 0x3813C000, 
				0x38140000, 0x38144000, 0x38148000, 0x3814C000, 0x38150000, 0x38154000, 0x38158000, 0x3815C000, 0x38160000, 0x38164000, 0x38168000, 0x3816C000, 0x38170000, 0x38174000, 0x38178000, 0x3817C000, 
				0x38180000, 0x38184000, 0x38188000, 0x3818C000, 0x38190000, 0x38194000, 0x38198000, 0x3819C000, 0x381A0000, 0x381A4000, 0x381A8000, 0x381AC000, 0x381B0000, 0x381B4000, 0x381B8000, 0x381BC000, 
				0x381C0000, 0x381C4000, 0x381C8000, 0x381CC000, 0x381D0000, 0x381D4000, 0x381D8000, 0x381DC000, 0x381E0000, 0x381E4000, 0x381E8000, 0x381EC000, 0x381F0000, 0x381F4000, 0x381F8000, 0x381FC000, 
				0x38200000, 0x38204000, 0x38208000, 0x3820C000, 0x38210000, 0x38214000, 0x38218000, 0x3821C000, 0x38220000, 0x38224000, 0x38228000, 0x3822C000, 0x38230000, 0x38234000, 0x38238000, 0x3823C000, 
				0x38240000, 0x38244000, 0x38248000, 0x3824C000, 0x38250000, 0x38254000, 0x38258000, 0x3825C000, 0x38260000, 0x38264000, 0x38268000, 0x3826C000, 0x38270000, 0x38274000, 0x38278000, 0x3827C000, 
				0x38280000, 0x38284000, 0x38288000, 0x3828C000, 0x38290000, 0x38294000, 0x38298000, 0x3829C000, 0x382A0000, 0x382A4000, 0x382A8000, 0x382AC000, 0x382B0000, 0x382B4000, 0x382B8000, 0x382BC000, 
				0x382C0000, 0x382C4000, 0x382C8000, 0x382CC000, 0x382D0000, 0x382D4000, 0x382D8000, 0x382DC000, 0x382E0000, 0x382E4000, 0x382E8000, 0x382EC000, 0x382F0000, 0x382F4000, 0x382F8000, 0x382FC000, 
				0x38300000, 0x38304000, 0x38308000, 0x3830C000, 0x38310000, 0x38314000, 0x38318000, 0x3831C000, 0x38320000, 0x38324000, 0x38328000, 0x3832C000, 0x38330000, 0x38334000, 0x38338000, 0x3833C000, 
				0x38340000, 0x38344000, 0x38348000, 0x3834C000, 0x38350000, 0x38354000, 0x38358000, 0x3835C000, 0x38360000, 0x38364000, 0x38368000, 0x3836C000, 0x38370000, 0x38374000, 0x38378000, 0x3837C000, 
				0x38380000, 0x38384000, 0x38388000, 0x3838C000, 0x38390000, 0x38394000, 0x38398000, 0x3839C000, 0x383A0000, 0x383A4000, 0x383A8000, 0x383AC000, 0x383B0000, 0x383B4000, 0x383B8000, 0x383BC000, 
				0x383C0000, 0x383C4000, 0x383C8000, 0x383CC000, 0x383D0000, 0x383D4000, 0x383D8000, 0x383DC000, 0x383E0000, 0x383E4000, 0x383E8000, 0x383EC000, 0x383F0000, 0x383F4000, 0x383F8000, 0x383FC000, 
				0x38400000, 0x38404000, 0x38408000, 0x3840C000, 0x38410000, 0x38414000, 0x38418000, 0x3841C000, 0x38420000, 0x38424000, 0x38428000, 0x3842C000, 0x38430000, 0x38434000, 0x38438000, 0x3843C000, 
				0x38440000, 0x38444000, 0x38448000, 0x3844C000, 0x38450000, 0x38454000, 0x38458000, 0x3845C000, 0x38460000, 0x38464000, 0x38468000, 0x3846C000, 0x38470000, 0x38474000, 0x38478000, 0x3847C000, 
				0x38480000, 0x38484000, 0x38488000, 0x3848C000, 0x38490000, 0x38494000, 0x38498000, 0x3849C000, 0x384A0000, 0x384A4000, 0x384A8000, 0x384AC000, 0x384B0000, 0x384B4000, 0x384B8000, 0x384BC000, 
				0x384C0000, 0x384C4000, 0x384C8000, 0x384CC000, 0x384D0000, 0x384D4000, 0x384D8000, 0x384DC000, 0x384E0000, 0x384E4000, 0x384E8000, 0x384EC000, 0x384F0000, 0x384F4000, 0x384F8000, 0x384FC000, 
				0x38500000, 0x38504000, 0x38508000, 0x3850C000, 0x38510000, 0x38514000, 0x38518000, 0x3851C000, 0x38520000, 0x38524000, 0x38528000, 0x3852C000, 0x38530000, 0x38534000, 0x38538000, 0x3853C000, 
				0x38540000, 0x38544000, 0x38548000, 0x3854C000, 0x38550000, 0x38554000, 0x38558000, 0x3855C000, 0x38560000, 0x38564000, 0x38568000, 0x3856C000, 0x38570000, 0x38574000, 0x38578000, 0x3857C000, 
				0x38580000, 0x38584000, 0x38588000, 0x3858C000, 0x38590000, 0x38594000, 0x38598000, 0x3859C000, 0x385A0000, 0x385A4000, 0x385A8000, 0x385AC000, 0x385B0000, 0x385B4000, 0x385B8000, 0x385BC000, 
				0x385C0000, 0x385C4000, 0x385C8000, 0x385CC000, 0x385D0000, 0x385D4000, 0x385D8000, 0x385DC000, 0x385E0000, 0x385E4000, 0x385E8000, 0x385EC000, 0x385F0000, 0x385F4000, 0x385F8000, 0x385FC000, 
				0x38600000, 0x38604000, 0x38608000, 0x3860C000, 0x38610000, 0x38614000, 0x38618000, 0x3861C000, 0x38620000, 0x38624000, 0x38628000, 0x3862C000, 0x38630000, 0x38634000, 0x38638000, 0x3863C000, 
				0x38640000, 0x38644000, 0x38648000, 0x3864C000, 0x38650000, 0x38654000, 0x38658000, 0x3865C000, 0x38660000, 0x38664000, 0x38668000, 0x3866C000, 0x38670000, 0x38674000, 0x38678000, 0x3867C000, 
				0x38680000, 0x38684000, 0x38688000, 0x3868C000, 0x38690000, 0x38694000, 0x38698000, 0x3869C000, 0x386A0000, 0x386A4000, 0x386A8000, 0x386AC000, 0x386B0000, 0x386B4000, 0x386B8000, 0x386BC000, 
				0x386C0000, 0x386C4000, 0x386C8000, 0x386CC000, 0x386D0000, 0x386D4000, 0x386D8000, 0x386DC000, 0x386E0000, 0x386E4000, 0x386E8000, 0x386EC000, 0x386F0000, 0x386F4000, 0x386F8000, 0x386FC000, 
				0x38700000, 0x38704000, 0x38708000, 0x3870C000, 0x38710000, 0x38714000, 0x38718000, 0x3871C000, 0x38720000, 0x38724000, 0x38728000, 0x3872C000, 0x38730000, 0x38734000, 0x38738000, 0x3873C000, 
				0x38740000, 0x38744000, 0x38748000, 0x3874C000, 0x38750000, 0x38754000, 0x38758000, 0x3875C000, 0x38760000, 0x38764000, 0x38768000, 0x3876C000, 0x38770000, 0x38774000, 0x38778000, 0x3877C000, 
				0x38780000, 0x38784000, 0x38788000, 0x3878C000, 0x38790000, 0x38794000, 0x38798000, 0x3879C000, 0x387A0000, 0x387A4000, 0x387A8000, 0x387AC000, 0x387B0000, 0x387B4000, 0x387B8000, 0x387BC000, 
				0x387C0000, 0x387C4000, 0x387C8000, 0x387CC000, 0x387D0000, 0x387D4000, 0x387D8000, 0x387DC000, 0x387E0000, 0x387E4000, 0x387E8000, 0x387EC000, 0x387F0000, 0x387F4000, 0x387F8000, 0x387FC000, 
				0x38000000, 0x38002000, 0x38004000, 0x38006000, 0x38008000, 0x3800A000, 0x3800C000, 0x3800E000, 0x38010000, 0x38012000, 0x38014000, 0x38016000, 0x38018000, 0x3801A000, 0x3801C000, 0x3801E000, 
				0x38020000, 0x38022000, 0x38024000, 0x38026000, 0x38028000, 0x3802A000, 0x3802C000, 0x3802E000, 0x38030000, 0x38032000, 0x38034000, 0x38036000, 0x38038000, 0x3803A000, 0x3803C000, 0x3803E000, 
				0x38040000, 0x38042000, 0x38044000, 0x38046000, 0x38048000, 0x3804A000, 0x3804C000, 0x3804E000, 0x38050000, 0x38052000, 0x38054000, 0x38056000, 0x38058000, 0x3805A000, 0x3805C000, 0x3805E000, 
				0x38060000, 0x38062000, 0x38064000, 0x38066000, 0x38068000, 0x3806A000, 0x3806C000, 0x3806E000, 0x38070000, 0x38072000, 0x38074000, 0x38076000, 0x38078000, 0x3807A000, 0x3807C000, 0x3807E000, 
				0x38080000, 0x38082000, 0x38084000, 0x38086000, 0x38088000, 0x3808A000, 0x3808C000, 0x3808E000, 0x38090000, 0x38092000, 0x38094000, 0x38096000, 0x38098000, 0x3809A000, 0x3809C000, 0x3809E000, 
				0x380A0000, 0x380A2000, 0x380A4000, 0x380A6000, 0x380A8000, 0x380AA000, 0x380AC000, 0x380AE000, 0x380B0000, 0x380B2000, 0x380B4000, 0x380B6000, 0x380B8000, 0x380BA000, 0x380BC000, 0x380BE000, 
				0x380C0000, 0x380C2000, 0x380C4000, 0x380C6000, 0x380C8000, 0x380CA000, 0x380CC000, 0x380CE000, 0x380D0000, 0x380D2000, 0x380D4000, 0x380D6000, 0x380D8000, 0x380DA000, 0x380DC000, 0x380DE000, 
				0x380E0000, 0x380E2000, 0x380E4000, 0x380E6000, 0x380E8000, 0x380EA000, 0x380EC000, 0x380EE000, 0x380F0000, 0x380F2000, 0x380F4000, 0x380F6000, 0x380F8000, 0x380FA000, 0x380FC000, 0x380FE000, 
				0x38100000, 0x38102000, 0x38104000, 0x38106000, 0x38108000, 0x3810A000, 0x3810C000, 0x3810E000, 0x38110000, 0x38112000, 0x38114000, 0x38116000, 0x38118000, 0x3811A000, 0x3811C000, 0x3811E000, 
				0x38120000, 0x38122000, 0x38124000, 0x38126000, 0x38128000, 0x3812A000, 0x3812C000, 0x3812E000, 0x38130000, 0x38132000, 0x38134000, 0x38136000, 0x38138000, 0x3813A000, 0x3813C000, 0x3813E000, 
				0x38140000, 0x38142000, 0x38144000, 0x38146000, 0x38148000, 0x3814A000, 0x3814C000, 0x3814E000, 0x38150000, 0x38152000, 0x38154000, 0x38156000, 0x38158000, 0x3815A000, 0x3815C000, 0x3815E000, 
				0x38160000, 0x38162000, 0x38164000, 0x38166000, 0x38168000, 0x3816A000, 0x3816C000, 0x3816E000, 0x38170000, 0x38172000, 0x38174000, 0x38176000, 0x38178000, 0x3817A000, 0x3817C000, 0x3817E000, 
				0x38180000, 0x38182000, 0x38184000, 0x38186000, 0x38188000, 0x3818A000, 0x3818C000, 0x3818E000, 0x38190000, 0x38192000, 0x38194000, 0x38196000, 0x38198000, 0x3819A000, 0x3819C000, 0x3819E000, 
				0x381A0000, 0x381A2000, 0x381A4000, 0x381A6000, 0x381A8000, 0x381AA000, 0x381AC000, 0x381AE000, 0x381B0000, 0x381B2000, 0x381B4000, 0x381B6000, 0x381B8000, 0x381BA000, 0x381BC000, 0x381BE000, 
				0x381C0000, 0x381C2000, 0x381C4000, 0x381C6000, 0x381C8000, 0x381CA000, 0x381CC000, 0x381CE000, 0x381D0000, 0x381D2000, 0x381D4000, 0x381D6000, 0x381D8000, 0x381DA000, 0x381DC000, 0x381DE000, 
				0x381E0000, 0x381E2000, 0x381E4000, 0x381E6000, 0x381E8000, 0x381EA000, 0x381EC000, 0x381EE000, 0x381F0000, 0x381F2000, 0x381F4000, 0x381F6000, 0x381F8000, 0x381FA000, 0x381FC000, 0x381FE000, 
				0x38200000, 0x38202000, 0x38204000, 0x38206000, 0x38208000, 0x3820A000, 0x3820C000, 0x3820E000, 0x38210000, 0x38212000, 0x38214000, 0x38216000, 0x38218000, 0x3821A000, 0x3821C000, 0x3821E000, 
				0x38220000, 0x38222000, 0x38224000, 0x38226000, 0x38228000, 0x3822A000, 0x3822C000, 0x3822E000, 0x38230000, 0x38232000, 0x38234000, 0x38236000, 0x38238000, 0x3823A000, 0x3823C000, 0x3823E000, 
				0x38240000, 0x38242000, 0x38244000, 0x38246000, 0x38248000, 0x3824A000, 0x3824C000, 0x3824E000, 0x38250000, 0x38252000, 0x38254000, 0x38256000, 0x38258000, 0x3825A000, 0x3825C000, 0x3825E000, 
				0x38260000, 0x38262000, 0x38264000, 0x38266000, 0x38268000, 0x3826A000, 0x3826C000, 0x3826E000, 0x38270000, 0x38272000, 0x38274000, 0x38276000, 0x38278000, 0x3827A000, 0x3827C000, 0x3827E000, 
				0x38280000, 0x38282000, 0x38284000, 0x38286000, 0x38288000, 0x3828A000, 0x3828C000, 0x3828E000, 0x38290000, 0x38292000, 0x38294000, 0x38296000, 0x38298000, 0x3829A000, 0x3829C000, 0x3829E000, 
				0x382A0000, 0x382A2000, 0x382A4000, 0x382A6000, 0x382A8000, 0x382AA000, 0x382AC000, 0x382AE000, 0x382B0000, 0x382B2000, 0x382B4000, 0x382B6000, 0x382B8000, 0x382BA000, 0x382BC000, 0x382BE000, 
				0x382C0000, 0x382C2000, 0x382C4000, 0x382C6000, 0x382C8000, 0x382CA000, 0x382CC000, 0x382CE000, 0x382D0000, 0x382D2000, 0x382D4000, 0x382D6000, 0x382D8000, 0x382DA000, 0x382DC000, 0x382DE000, 
				0x382E0000, 0x382E2000, 0x382E4000, 0x382E6000, 0x382E8000, 0x382EA000, 0x382EC000, 0x382EE000, 0x382F0000, 0x382F2000, 0x382F4000, 0x382F6000, 0x382F8000, 0x382FA000, 0x382FC000, 0x382FE000, 
				0x38300000, 0x38302000, 0x38304000, 0x38306000, 0x38308000, 0x3830A000, 0x3830C000, 0x3830E000, 0x38310000, 0x38312000, 0x38314000, 0x38316000, 0x38318000, 0x3831A000, 0x3831C000, 0x3831E000, 
				0x38320000, 0x38322000, 0x38324000, 0x38326000, 0x38328000, 0x3832A000, 0x3832C000, 0x3832E000, 0x38330000, 0x38332000, 0x38334000, 0x38336000, 0x38338000, 0x3833A000, 0x3833C000, 0x3833E000, 
				0x38340000, 0x38342000, 0x38344000, 0x38346000, 0x38348000, 0x3834A000, 0x3834C000, 0x3834E000, 0x38350000, 0x38352000, 0x38354000, 0x38356000, 0x38358000, 0x3835A000, 0x3835C000, 0x3835E000, 
				0x38360000, 0x38362000, 0x38364000, 0x38366000, 0x38368000, 0x3836A000, 0x3836C000, 0x3836E000, 0x38370000, 0x38372000, 0x38374000, 0x38376000, 0x38378000, 0x3837A000, 0x3837C000, 0x3837E000, 
				0x38380000, 0x38382000, 0x38384000, 0x38386000, 0x38388000, 0x3838A000, 0x3838C000, 0x3838E000, 0x38390000, 0x38392000, 0x38394000, 0x38396000, 0x38398000, 0x3839A000, 0x3839C000, 0x3839E000, 
				0x383A0000, 0x383A2000, 0x383A4000, 0x383A6000, 0x383A8000, 0x383AA000, 0x383AC000, 0x383AE000, 0x383B0000, 0x383B2000, 0x383B4000, 0x383B6000, 0x383B8000, 0x383BA000, 0x383BC000, 0x383BE000, 
				0x383C0000, 0x383C2000, 0x383C4000, 0x383C6000, 0x383C8000, 0x383CA000, 0x383CC000, 0x383CE000, 0x383D0000, 0x383D2000, 0x383D4000, 0x383D6000, 0x383D8000, 0x383DA000, 0x383DC000, 0x383DE000, 
				0x383E0000, 0x383E2000, 0x383E4000, 0x383E6000, 0x383E8000, 0x383EA000, 0x383EC000, 0x383EE000, 0x383F0000, 0x383F2000, 0x383F4000, 0x383F6000, 0x383F8000, 0x383FA000, 0x383FC000, 0x383FE000, 
				0x38400000, 0x38402000, 0x38404000, 0x38406000, 0x38408000, 0x3840A000, 0x3840C000, 0x3840E000, 0x38410000, 0x38412000, 0x38414000, 0x38416000, 0x38418000, 0x3841A000, 0x3841C000, 0x3841E000, 
				0x38420000, 0x38422000, 0x38424000, 0x38426000, 0x38428000, 0x3842A000, 0x3842C000, 0x3842E000, 0x38430000, 0x38432000, 0x38434000, 0x38436000, 0x38438000, 0x3843A000, 0x3843C000, 0x3843E000, 
				0x38440000, 0x38442000, 0x38444000, 0x38446000, 0x38448000, 0x3844A000, 0x3844C000, 0x3844E000, 0x38450000, 0x38452000, 0x38454000, 0x38456000, 0x38458000, 0x3845A000, 0x3845C000, 0x3845E000, 
				0x38460000, 0x38462000, 0x38464000, 0x38466000, 0x38468000, 0x3846A000, 0x3846C000, 0x3846E000, 0x38470000, 0x38472000, 0x38474000, 0x38476000, 0x38478000, 0x3847A000, 0x3847C000, 0x3847E000, 
				0x38480000, 0x38482000, 0x38484000, 0x38486000, 0x38488000, 0x3848A000, 0x3848C000, 0x3848E000, 0x38490000, 0x38492000, 0x38494000, 0x38496000, 0x38498000, 0x3849A000, 0x3849C000, 0x3849E000, 
				0x384A0000, 0x384A2000, 0x384A4000, 0x384A6000, 0x384A8000, 0x384AA000, 0x384AC000, 0x384AE000, 0x384B0000, 0x384B2000, 0x384B4000, 0x384B6000, 0x384B8000, 0x384BA000, 0x384BC000, 0x384BE000, 
				0x384C0000, 0x384C2000, 0x384C4000, 0x384C6000, 0x384C8000, 0x384CA000, 0x384CC000, 0x384CE000, 0x384D0000, 0x384D2000, 0x384D4000, 0x384D6000, 0x384D8000, 0x384DA000, 0x384DC000, 0x384DE000, 
				0x384E0000, 0x384E2000, 0x384E4000, 0x384E6000, 0x384E8000, 0x384EA000, 0x384EC000, 0x384EE000, 0x384F0000, 0x384F2000, 0x384F4000, 0x384F6000, 0x384F8000, 0x384FA000, 0x384FC000, 0x384FE000, 
				0x38500000, 0x38502000, 0x38504000, 0x38506000, 0x38508000, 0x3850A000, 0x3850C000, 0x3850E000, 0x38510000, 0x38512000, 0x38514000, 0x38516000, 0x38518000, 0x3851A000, 0x3851C000, 0x3851E000, 
				0x38520000, 0x38522000, 0x38524000, 0x38526000, 0x38528000, 0x3852A000, 0x3852C000, 0x3852E000, 0x38530000, 0x38532000, 0x38534000, 0x38536000, 0x38538000, 0x3853A000, 0x3853C000, 0x3853E000, 
				0x38540000, 0x38542000, 0x38544000, 0x38546000, 0x38548000, 0x3854A000, 0x3854C000, 0x3854E000, 0x38550000, 0x38552000, 0x38554000, 0x38556000, 0x38558000, 0x3855A000, 0x3855C000, 0x3855E000, 
				0x38560000, 0x38562000, 0x38564000, 0x38566000, 0x38568000, 0x3856A000, 0x3856C000, 0x3856E000, 0x38570000, 0x38572000, 0x38574000, 0x38576000, 0x38578000, 0x3857A000, 0x3857C000, 0x3857E000, 
				0x38580000, 0x38582000, 0x38584000, 0x38586000, 0x38588000, 0x3858A000, 0x3858C000, 0x3858E000, 0x38590000, 0x38592000, 0x38594000, 0x38596000, 0x38598000, 0x3859A000, 0x3859C000, 0x3859E000, 
				0x385A0000, 0x385A2000, 0x385A4000, 0x385A6000, 0x385A8000, 0x385AA000, 0x385AC000, 0x385AE000, 0x385B0000, 0x385B2000, 0x385B4000, 0x385B6000, 0x385B8000, 0x385BA000, 0x385BC000, 0x385BE000, 
				0x385C0000, 0x385C2000, 0x385C4000, 0x385C6000, 0x385C8000, 0x385CA000, 0x385CC000, 0x385CE000, 0x385D0000, 0x385D2000, 0x385D4000, 0x385D6000, 0x385D8000, 0x385DA000, 0x385DC000, 0x385DE000, 
				0x385E0000, 0x385E2000, 0x385E4000, 0x385E6000, 0x385E8000, 0x385EA000, 0x385EC000, 0x385EE000, 0x385F0000, 0x385F2000, 0x385F4000, 0x385F6000, 0x385F8000, 0x385FA000, 0x385FC000, 0x385FE000, 
				0x38600000, 0x38602000, 0x38604000, 0x38606000, 0x38608000, 0x3860A000, 0x3860C000, 0x3860E000, 0x38610000, 0x38612000, 0x38614000, 0x38616000, 0x38618000, 0x3861A000, 0x3861C000, 0x3861E000, 
				0x38620000, 0x38622000, 0x38624000, 0x38626000, 0x38628000, 0x3862A000, 0x3862C000, 0x3862E000, 0x38630000, 0x38632000, 0x38634000, 0x38636000, 0x38638000, 0x3863A000, 0x3863C000, 0x3863E000, 
				0x38640000, 0x38642000, 0x38644000, 0x38646000, 0x38648000, 0x3864A000, 0x3864C000, 0x3864E000, 0x38650000, 0x38652000, 0x38654000, 0x38656000, 0x38658000, 0x3865A000, 0x3865C000, 0x3865E000, 
				0x38660000, 0x38662000, 0x38664000, 0x38666000, 0x38668000, 0x3866A000, 0x3866C000, 0x3866E000, 0x38670000, 0x38672000, 0x38674000, 0x38676000, 0x38678000, 0x3867A000, 0x3867C000, 0x3867E000, 
				0x38680000, 0x38682000, 0x38684000, 0x38686000, 0x38688000, 0x3868A000, 0x3868C000, 0x3868E000, 0x38690000, 0x38692000, 0x38694000, 0x38696000, 0x38698000, 0x3869A000, 0x3869C000, 0x3869E000, 
				0x386A0000, 0x386A2000, 0x386A4000, 0x386A6000, 0x386A8000, 0x386AA000, 0x386AC000, 0x386AE000, 0x386B0000, 0x386B2000, 0x386B4000, 0x386B6000, 0x386B8000, 0x386BA000, 0x386BC000, 0x386BE000, 
				0x386C0000, 0x386C2000, 0x386C4000, 0x386C6000, 0x386C8000, 0x386CA000, 0x386CC000, 0x386CE000, 0x386D0000, 0x386D2000, 0x386D4000, 0x386D6000, 0x386D8000, 0x386DA000, 0x386DC000, 0x386DE000, 
				0x386E0000, 0x386E2000, 0x386E4000, 0x386E6000, 0x386E8000, 0x386EA000, 0x386EC000, 0x386EE000, 0x386F0000, 0x386F2000, 0x386F4000, 0x386F6000, 0x386F8000, 0x386FA000, 0x386FC000, 0x386FE000, 
				0x38700000, 0x38702000, 0x38704000, 0x38706000, 0x38708000, 0x3870A000, 0x3870C000, 0x3870E000, 0x38710000, 0x38712000, 0x38714000, 0x38716000, 0x38718000, 0x3871A000, 0x3871C000, 0x3871E000, 
				0x38720000, 0x38722000, 0x38724000, 0x38726000, 0x38728000, 0x3872A000, 0x3872C000, 0x3872E000, 0x38730000, 0x38732000, 0x38734000, 0x38736000, 0x38738000, 0x3873A000, 0x3873C000, 0x3873E000, 
				0x38740000, 0x38742000, 0x38744000, 0x38746000, 0x38748000, 0x3874A000, 0x3874C000, 0x3874E000, 0x38750000, 0x38752000, 0x38754000, 0x38756000, 0x38758000, 0x3875A000, 0x3875C000, 0x3875E000, 
				0x38760000, 0x38762000, 0x38764000, 0x38766000, 0x38768000, 0x3876A000, 0x3876C000, 0x3876E000, 0x38770000, 0x38772000, 0x38774000, 0x38776000, 0x38778000, 0x3877A000, 0x3877C000, 0x3877E000, 
				0x38780000, 0x38782000, 0x38784000, 0x38786000, 0x38788000, 0x3878A000, 0x3878C000, 0x3878E000, 0x38790000, 0x38792000, 0x38794000, 0x38796000, 0x38798000, 0x3879A000, 0x3879C000, 0x3879E000, 
				0x387A0000, 0x387A2000, 0x387A4000, 0x387A6000, 0x387A8000, 0x387AA000, 0x387AC000, 0x387AE000, 0x387B0000, 0x387B2000, 0x387B4000, 0x387B6000, 0x387B8000, 0x387BA000, 0x387BC000, 0x387BE000, 
				0x387C0000, 0x387C2000, 0x387C4000, 0x387C6000, 0x387C8000, 0x387CA000, 0x387CC000, 0x387CE000, 0x387D0000, 0x387D2000, 0x387D4000, 0x387D6000, 0x387D8000, 0x387DA000, 0x387DC000, 0x387DE000, 
				0x387E0000, 0x387E2000, 0x387E4000, 0x387E6000, 0x387E8000, 0x387EA000, 0x387EC000, 0x387EE000, 0x387F0000, 0x387F2000, 0x387F4000, 0x387F6000, 0x387F8000, 0x387FA000, 0x387FC000, 0x387FE000 };
			static const bits_t<float> exponent_table[64] = {
				0x00000000, 0x00800000, 0x01000000, 0x01800000, 0x02000000, 0x02800000, 0x03000000, 0x03800000, 0x04000000, 0x04800000, 0x05000000, 0x05800000, 0x06000000, 0x06800000, 0x07000000, 0x07800000, 
				0x08000000, 0x08800000, 0x09000000, 0x09800000, 0x0A000000, 0x0A800000, 0x0B000000, 0x0B800000, 0x0C000000, 0x0C800000, 0x0D000000, 0x0D800000, 0x0E000000, 0x0E800000, 0x0F000000, 0x47800000, 
				0x80000000, 0x80800000, 0x81000000, 0x81800000, 0x82000000, 0x82800000, 0x83000000, 0x83800000, 0x84000000, 0x84800000, 0x85000000, 0x85800000, 0x86000000, 0x86800000, 0x87000000, 0x87800000, 
				0x88000000, 0x88800000, 0x89000000, 0x89800000, 0x8A000000, 0x8A800000, 0x8B000000, 0x8B800000, 0x8C000000, 0x8C800000, 0x8D000000, 0x8D800000, 0x8E000000, 0x8E800000, 0x8F000000, 0xC7800000 };
			static const unsigned short offset_table[64] = {
				0, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 
				0, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024, 1024 };
			bits_t<float> fbits = mantissa_table[offset_table[value>>10]+(value&0x3FF)] + exponent_table[value>>10];
		#endif
			float out;
			std::memcpy(&out, &fbits, sizeof(float));
			return out;
		#endif
		}

		/// Convert half-precision to IEEE double-precision.
		/// \param value half-precision value to convert
		/// \return double-precision value
		inline double half2float_impl(unsigned int value, double, true_type) {
		#if HALF_ENABLE_F16C_INTRINSICS
			return _mm_cvtsd_f64(_mm_cvtps_pd(_mm_cvtph_ps(_mm_cvtsi32_si128(value))));
		#else
			uint32 hi = static_cast<uint32>(value&0x8000) << 16;
			unsigned int abs = value & 0x7FFF;
			if(abs) {
				hi |= 0x3F000000 << static_cast<unsigned>(abs>=0x7C00);
				for(; abs<0x400; abs<<=1,hi-=0x100000) ;
				hi += static_cast<uint32>(abs) << 10;
			}
			bits_t<double> dbits = static_cast<bits_t<double>>(hi) << 32;
			double out;
			std::memcpy(&out, &dbits, sizeof(double));
			return out;
		#endif
		}

		/// Convert half-precision to non-IEEE floating-point.
		/// \tparam T type to convert to (builtin integer type)
		/// \param value half-precision value to convert
		/// \return floating-point value
		template<class T> T half2float_impl(unsigned int value, T, ...) {
			T out;
			unsigned int abs = value & 0x7FFF;
			if(abs > 0x7C00)
				out = (std::numeric_limits<T>::has_signaling_NaN && !(abs&0x200)) ? std::numeric_limits<T>::signaling_NaN() :
					std::numeric_limits<T>::has_quiet_NaN ? std::numeric_limits<T>::quiet_NaN() : T();
			else if(abs == 0x7C00)
				out = std::numeric_limits<T>::has_infinity ? std::numeric_limits<T>::infinity() : std::numeric_limits<T>::max();
			else if(abs > 0x3FF)
				out = std::ldexp(static_cast<T>((abs&0x3FF)|0x400), (abs>>10)-25);
			else
				out = std::ldexp(static_cast<T>(abs), -24);
			return (value&0x8000) ? -out : out;
		}

		/// Convert half-precision to floating-point.
		/// \tparam T type to convert to (builtin integer type)
		/// \param value half-precision value to convert
		/// \return floating-point value
		template<class T> T half2float(unsigned int value) {
			return half2float_impl(value, T(), bool_type<std::numeric_limits<T>::is_iec559&&sizeof(bits_t<T>)==sizeof(T)>());
		}

		/// Convert half-precision floating-point to integer.
		/// \tparam R rounding mode to use
		/// \tparam E `true` for round to even, `false` for round away from zero
		/// \tparam I `true` to raise INEXACT exception (if inexact), `false` to never raise it
		/// \tparam T type to convert to (buitlin integer type with at least 16 bits precision, excluding any implicit sign bits)
		/// \param value half-precision value to convert
		/// \return rounded integer value
		/// \exception FE_INVALID if value is not representable in type \a T
		/// \exception FE_INEXACT if value had to be rounded and \a I is `true`
		template<std::float_round_style R,bool E,bool I,class T> T half2int(unsigned int value) {
			unsigned int abs = value & 0x7FFF;
			if(abs >= 0x7C00) {
				raise(FE_INVALID);
				return (value&0x8000) ? std::numeric_limits<T>::min() : std::numeric_limits<T>::max();
			}
			if(abs < 0x3800) {
				raise(FE_INEXACT, I);
				return	(R==std::round_toward_infinity) ? T(~(value>>15)&(abs!=0)) :
						(R==std::round_toward_neg_infinity) ? -T(value>0x8000) :
						T();
			}
			int exp = 25 - (abs>>10);
			unsigned int m = (value&0x3FF) | 0x400;
			int32 i = static_cast<int32>((exp<=0) ? (m<<-exp) : ((m+(
				(R==std::round_to_nearest) ? ((1<<(exp-1))-(~(m>>exp)&E)) :
				(R==std::round_toward_infinity) ? (((1<<exp)-1)&((value>>15)-1)) :
				(R==std::round_toward_neg_infinity) ? (((1<<exp)-1)&-(value>>15)) : 0))>>exp));
			if((!std::numeric_limits<T>::is_signed && (value&0x8000)) || (std::numeric_limits<T>::digits<16 &&
				((value&0x8000) ? (-i<std::numeric_limits<T>::min()) : (i>std::numeric_limits<T>::max()))))
				raise(FE_INVALID);
			else if(I && exp > 0 && (m&((1<<exp)-1)))
				raise(FE_INEXACT);
			return static_cast<T>((value&0x8000) ? -i : i);
		}

		/// \}
		/// \name Mathematics
		/// \{

		/// upper part of 64-bit multiplication.
		/// \tparam R rounding mode to use
		/// \param x first factor
		/// \param y second factor
		/// \return upper 32 bit of \a x * \a y
		template<std::float_round_style R> uint32 mulhi(uint32 x, uint32 y) {
			uint32 xy = (x>>16) * (y&0xFFFF), yx = (x&0xFFFF) * (y>>16), c = (xy&0xFFFF) + (yx&0xFFFF) + (((x&0xFFFF)*(y&0xFFFF))>>16);
			return (x>>16)*(y>>16) + (xy>>16) + (yx>>16) + (c>>16) +
				((R==std::round_to_nearest) ? ((c>>15)&1) : (R==std::round_toward_infinity) ? ((c&0xFFFF)!=0) : 0);
		}

		/// 64-bit multiplication.
		/// \param x first factor
		/// \param y second factor
		/// \return upper 32 bit of \a x * \a y rounded to nearest
		inline uint32 multiply64(uint32 x, uint32 y) {
			return static_cast<uint32>((static_cast<unsigned long long>(x)*static_cast<unsigned long long>(y)+0x80000000)>>32);
		}

		/// 64-bit division.
		/// \param x upper 32 bit of dividend
		/// \param y divisor
		/// \param s variable to store sticky bit for rounding
		/// \return (\a x << 32) / \a y
		inline uint32 divide64(uint32 x, uint32 y, int &s) {
			unsigned long long xx = static_cast<unsigned long long>(x) << 32;
			return s = (xx%y!=0), static_cast<uint32>(xx/y);
		}

		/// Half precision positive modulus.
		/// \tparam Q `true` to compute full quotient, `false` else
		/// \tparam R `true` to compute signed remainder, `false` for positive remainder
		/// \param x first operand as positive finite half-precision value
		/// \param y second operand as positive finite half-precision value
		/// \param quo adress to store quotient at, `nullptr` if \a Q `false`
		/// \return modulus of \a x / \a y
		template<bool Q,bool R> unsigned int mod(unsigned int x, unsigned int y, int *quo = NULL) {
			unsigned int q = 0;
			if(x > y) {
				int absx = x, absy = y, expx = 0, expy = 0;
				for(; absx<0x400; absx<<=1,--expx) ;
				for(; absy<0x400; absy<<=1,--expy) ;
				expx += absx >> 10;
				expy += absy >> 10;
				int mx = (absx&0x3FF) | 0x400, my = (absy&0x3FF) | 0x400;
				for(int d=expx-expy; d; --d) {
					if(!Q && mx == my)
						return 0;
					if(mx >= my) {
						mx -= my;
						q += Q;
					}
					mx <<= 1;
					q <<= static_cast<int>(Q);
				}
				if(!Q && mx == my)
					return 0;
				if(mx >= my) {
					mx -= my;
					++q;
				}
				if(Q) {
					q &= (1<<(std::numeric_limits<int>::digits-1)) - 1;
					if(!mx)
						return *quo = q, 0;
				}
				for(; mx<0x400; mx<<=1,--expy) ;
				x = (expy>0) ? ((expy<<10)|(mx&0x3FF)) : (mx>>(1-expy));
			}
			if(R) {
				unsigned int a, b;
				if(y < 0x800) {
					a = (x<0x400) ? (x<<1) : (x+0x400);
					b = y;
				} else {
					a = x;
					b = y - 0x400;
				}
				if(a > b || (a == b && (q&1))) {
					int exp = (y>>10) + (y<=0x3FF), d = exp - (x>>10) - (x<=0x3FF);
					int m = (((y&0x3FF)|((y>0x3FF)<<10))<<1) - (((x&0x3FF)|((x>0x3FF)<<10))<<(1-d));
					for(; m<0x800 && exp>1; m<<=1,--exp) ;
					x = 0x8000 + ((exp-1)<<10) + (m>>1);
					q += Q;
				}
			}
			if(Q)
				*quo = q;
			return x;
		}

		/// Fixed point square root.
		/// \tparam F number of fractional bits
		/// \param r radicand in Q1.F fixed point format
		/// \param exp exponent
		/// \return square root as Q1.F/2
		template<unsigned int F> uint32 sqrt(uint32 &r, int &exp) {
			int i = exp & 1;
			r <<= i;
			exp = (exp-i) / 2;
			uint32 m = 0;
			for(uint32 bit=static_cast<uint32>(1)<<F; bit; bit>>=2) {
				if(r < m+bit)
					m >>= 1;
				else {
					r -= m + bit;
					m = (m>>1) + bit;
				}
			}
			return m;
		}

		/// Fixed point binary exponential.
		/// This uses the BKM algorithm in E-mode.
		/// \param m exponent in [0,1) as Q0.31
		/// \param n number of iterations (at most 32)
		/// \return 2 ^ \a m as Q1.31
		inline uint32 exp2(uint32 m, unsigned int n = 32) {
			static const uint32 logs[] = {
				0x80000000, 0x4AE00D1D, 0x2934F098, 0x15C01A3A, 0x0B31FB7D, 0x05AEB4DD, 0x02DCF2D1, 0x016FE50B,
				0x00B84E23, 0x005C3E10, 0x002E24CA, 0x001713D6, 0x000B8A47, 0x0005C53B, 0x0002E2A3, 0x00017153,
				0x0000B8AA, 0x00005C55, 0x00002E2B, 0x00001715, 0x00000B8B, 0x000005C5, 0x000002E3, 0x00000171,
				0x000000B9, 0x0000005C, 0x0000002E, 0x00000017, 0x0000000C, 0x00000006, 0x00000003, 0x00000001 };
			if(!m)
				return 0x80000000;
			uint32 mx = 0x80000000, my = 0;
			for(unsigned int i=1; i<n; ++i) {
				uint32 mz = my + logs[i];
				if(mz <= m) {
					my = mz;
					mx += mx >> i;
				}
			}
			return mx;
		}

		/// Fixed point binary logarithm.
		/// This uses the BKM algorithm in L-mode.
		/// \param m mantissa in [1,2) as Q1.30
		/// \param n number of iterations (at most 32)
		/// \return log2(\a m) as Q0.31
		inline uint32 log2(uint32 m, unsigned int n = 32) {
			static const uint32 logs[] = {
				0x80000000, 0x4AE00D1D, 0x2934F098, 0x15C01A3A, 0x0B31FB7D, 0x05AEB4DD, 0x02DCF2D1, 0x016FE50B,
				0x00B84E23, 0x005C3E10, 0x002E24CA, 0x001713D6, 0x000B8A47, 0x0005C53B, 0x0002E2A3, 0x00017153,
				0x0000B8AA, 0x00005C55, 0x00002E2B, 0x00001715, 0x00000B8B, 0x000005C5, 0x000002E3, 0x00000171,
				0x000000B9, 0x0000005C, 0x0000002E, 0x00000017, 0x0000000C, 0x00000006, 0x00000003, 0x00000001 };
			if(m == 0x40000000)
				return 0;
			uint32 mx = 0x40000000, my = 0;
			for(unsigned int i=1; i<n; ++i) {
				uint32 mz = mx + (mx>>i);
				if(mz <= m) {
					mx = mz;
					my += logs[i];
				}
			}
			return my;
		}

		/// Fixed point sine and cosine.
		/// This uses the CORDIC algorithm in rotation mode.
		/// \param mz angle in [-pi/2,pi/2] as Q1.30
		/// \param n number of iterations (at most 31)
		/// \return sine and cosine of \a mz as Q1.30
		inline std::pair<uint32,uint32> sincos(uint32 mz, unsigned int n = 31) {
			static const uint32 angles[] = {
				0x3243F6A9, 0x1DAC6705, 0x0FADBAFD, 0x07F56EA7, 0x03FEAB77, 0x01FFD55C, 0x00FFFAAB, 0x007FFF55,
				0x003FFFEB, 0x001FFFFD, 0x00100000, 0x00080000, 0x00040000, 0x00020000, 0x00010000, 0x00008000,
				0x00004000, 0x00002000, 0x00001000, 0x00000800, 0x00000400, 0x00000200, 0x00000100, 0x00000080,
				0x00000040, 0x00000020, 0x00000010, 0x00000008, 0x00000004, 0x00000002, 0x00000001 };
			uint32 mx = 0x26DD3B6A, my = 0;
			for(unsigned int i=0; i<n; ++i) {
				uint32 sign = sign_mask(mz);
				uint32 tx = mx - (arithmetic_shift(my, i)^sign) + sign;
				uint32 ty = my + (arithmetic_shift(mx, i)^sign) - sign;
				mx = tx; my = ty; mz -= (angles[i]^sign) - sign;
			}
			return std::make_pair(my, mx);
		}

		/// Fixed point arc tangent.
		/// This uses the CORDIC algorithm in vectoring mode.
		/// \param my y coordinate as Q0.30
		/// \param mx x coordinate as Q0.30
		/// \param n number of iterations (at most 31)
		/// \return arc tangent of \a my / \a mx as Q1.30
		inline uint32 atan2(uint32 my, uint32 mx, unsigned int n = 31) {
			static const uint32 angles[] = {
				0x3243F6A9, 0x1DAC6705, 0x0FADBAFD, 0x07F56EA7, 0x03FEAB77, 0x01FFD55C, 0x00FFFAAB, 0x007FFF55,
				0x003FFFEB, 0x001FFFFD, 0x00100000, 0x00080000, 0x00040000, 0x00020000, 0x00010000, 0x00008000,
				0x00004000, 0x00002000, 0x00001000, 0x00000800, 0x00000400, 0x00000200, 0x00000100, 0x00000080,
				0x00000040, 0x00000020, 0x00000010, 0x00000008, 0x00000004, 0x00000002, 0x00000001 };
			uint32 mz = 0;
			for(unsigned int i=0; i<n; ++i) {
				uint32 sign = sign_mask(my);
				uint32 tx = mx + (arithmetic_shift(my, i)^sign) - sign;
				uint32 ty = my - (arithmetic_shift(mx, i)^sign) + sign;
				mx = tx; my = ty; mz += (angles[i]^sign) - sign;
			}
			return mz;
		}

		/// Reduce argument for trigonometric functions.
		/// \param abs half-precision floating-point value
		/// \param k value to take quarter period
		/// \return \a abs reduced to [-pi/4,pi/4] as Q0.30
		inline uint32 angle_arg(unsigned int abs, int &k) {
			uint32 m = (abs&0x3FF) | ((abs>0x3FF)<<10);
			int exp = (abs>>10) + (abs<=0x3FF) - 15;
			if(abs < 0x3A48)
				return k = 0, m << (exp+20);
			unsigned long long y = m * 0xA2F9836E4E442, mask = (1ULL<<(62-exp)) - 1, yi = (y+(mask>>1)) & ~mask, f = y - yi;
			uint32 sign = -static_cast<uint32>(f>>63);
			k = static_cast<int>(yi>>(62-exp));
			return (multiply64(static_cast<uint32>((sign ? -f : f)>>(31-exp)), 0xC90FDAA2)^sign) - sign;
		}

		/// Get arguments for atan2 function.
		/// \param abs half-precision floating-point value
		/// \return \a abs and sqrt(1 - \a abs^2) as Q0.30
		inline std::pair<uint32,uint32> atan2_args(unsigned int abs) {
			int exp = -15;
			for(; abs<0x400; abs<<=1,--exp) ;
			exp += abs >> 10;
			uint32 my = ((abs&0x3FF)|0x400) << 5, r = my * my;
			int rexp = 2 * exp;
			r = 0x40000000 - ((rexp>-31) ? ((r>>-rexp)|((r&((static_cast<uint32>(1)<<-rexp)-1))!=0)) : 1);
			for(rexp=0; r<0x40000000; r<<=1,--rexp) ;
			uint32 mx = sqrt<30>(r, rexp);
			int d = exp - rexp;
			if(d < 0)
				return std::make_pair((d<-14) ? ((my>>(-d-14))+((my>>(-d-15))&1)) : (my<<(14+d)), (mx<<14)+(r<<13)/mx);
			if(d > 0)
				return std::make_pair(my<<14, (d>14) ? ((mx>>(d-14))+((mx>>(d-15))&1)) : ((d==14) ? mx : ((mx<<(14-d))+(r<<(13-d))/mx)));
			return std::make_pair(my<<13, (mx<<13)+(r<<12)/mx);
		}

		/// Get exponentials for hyperbolic computation
		/// \param abs half-precision floating-point value
		/// \param exp variable to take unbiased exponent of larger result
		/// \param n number of BKM iterations (at most 32)
		/// \return exp(abs) and exp(-\a abs) as Q1.31 with same exponent
		inline std::pair<uint32,uint32> hyperbolic_args(unsigned int abs, int &exp, unsigned int n = 32) {
			uint32 mx = detail::multiply64(static_cast<uint32>((abs&0x3FF)+((abs>0x3FF)<<10))<<21, 0xB8AA3B29), my;
			int e = (abs>>10) + (abs<=0x3FF);
			if(e < 14) {
				exp = 0;
				mx >>= 14 - e;
			} else {
				exp = mx >> (45-e);
				mx = (mx<<(e-14)) & 0x7FFFFFFF;
			}
			mx = exp2(mx, n);
			int d = exp << 1, s;
			if(mx > 0x80000000) {
				my = divide64(0x80000000, mx, s);
				my |= s;
				++d;
			} else
				my = mx;
			return std::make_pair(mx, (d<31) ? ((my>>d)|((my&((static_cast<uint32>(1)<<d)-1))!=0)) : 1);
		}

		/// Postprocessing for binary exponential.
		/// \tparam R rounding mode to use
		/// \tparam I `true` to always raise INEXACT exception, `false` to raise only for rounded results
		/// \param m mantissa as Q1.31
		/// \param exp absolute value of unbiased exponent
		/// \param esign sign of actual exponent
		/// \param sign sign bit of result
		/// \return value converted to half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded or \a I is `true`
		template<std::float_round_style R,bool I> unsigned int exp2_post(uint32 m, int exp, bool esign, unsigned int sign = 0) {
			int s = 0;
			if(esign) {
				if(m > 0x80000000) {
					m = divide64(0x80000000, m, s);
					++exp;
				}
				if(exp > 25)
					return underflow<R>(sign);
				else if(exp == 25)
					return rounded<R,I>(sign, 1, (m&0x7FFFFFFF)!=0);
				exp = -exp;
			} else if(exp > 15)
				return overflow<R>(sign);
			return fixed2half<R,31,false,false,I>(m, exp+14, sign, s);
		}

		/// Postprocessing for binary logarithm.
		/// \tparam R rounding mode to use
		/// \tparam L logarithm for base transformation as Q1.31
		/// \param m fractional part of logarithm as Q0.31
		/// \param ilog signed integer part of logarithm
		/// \param exp biased exponent of result
		/// \param sign sign bit of result
		/// \return value base-transformed and converted to half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if no other exception occurred
		template<std::float_round_style R,uint32 L> unsigned int log2_post(uint32 m, int ilog, int exp, unsigned int sign = 0) {
			uint32 msign = sign_mask(ilog);
			m = (((static_cast<uint32>(ilog)<<27)+(m>>4))^msign) - msign;
			if(!m)
				return 0;
			for(; m<0x80000000; m<<=1,--exp) ;
			int i = m >= L, s;
			exp += i;
			m >>= 1 + i;
			sign ^= msign & 0x8000;
			if(exp < -11)
				return underflow<R>(sign);
			m = divide64(m, L, s);
			return fixed2half<R,30,false,false,true>(m, exp, sign, 1);
		}

		/// Hypotenuse square root and postprocessing.
		/// \tparam R rounding mode to use
		/// \param r mantissa as Q2.30
		/// \param exp unbiased exponent
		/// \return square root converted to half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if value had to be rounded
		template<std::float_round_style R> unsigned int hypot_post(uint32 r, int exp) {
			int i = r >> 31;
			if((exp+=i) > 46)
				return overflow<R>();
			if(exp < -34)
				return underflow<R>();
			r = (r>>i) | (r&i);
			uint32 m = sqrt<30>(r, exp+=15);
			return fixed2half<R,15,false,false,false>(m, exp-1, 0, r!=0);
		}

		/// Division and postprocessing for tangents.
		/// \tparam R rounding mode to use
		/// \param my dividend as Q1.31
		/// \param mx divisor as Q1.31
		/// \param exp biased exponent of result
		/// \param sign sign bit of result
		/// \return quotient converted to half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if no other exception occurred
		template<std::float_round_style R> unsigned int tangent_post(uint32 my, uint32 mx, int exp, unsigned int sign = 0) {
			int i = my >= mx, s;
			exp += i;
			if(exp > 29)
				return overflow<R>(sign);
			if(exp < -11)
				return underflow<R>(sign);
			uint32 m = divide64(my>>(i+1), mx, s);
			return fixed2half<R,30,false,false,true>(m, exp, sign, s);
		}

		/// Area function and postprocessing.
		/// This computes the value directly in Q2.30 using the representation `asinh|acosh(x) = log(x+sqrt(x^2+|-1))`.
		/// \tparam R rounding mode to use
		/// \tparam S `true` for asinh, `false` for acosh
		/// \param arg half-precision argument
		/// \return asinh|acosh(\a arg) converted to half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if no other exception occurred
		template<std::float_round_style R,bool S> unsigned int area(unsigned int arg) {
			int abs = arg & 0x7FFF, expx = (abs>>10) + (abs<=0x3FF) - 15, expy = -15, ilog, i;
			uint32 mx = static_cast<uint32>((abs&0x3FF)|((abs>0x3FF)<<10)) << 20, my, r;
			for(; abs<0x400; abs<<=1,--expy) ;
			expy += abs >> 10;
			r = ((abs&0x3FF)|0x400) << 5;
			r *= r;
			i = r >> 31;
			expy = 2*expy + i;
			r >>= i;
			if(S) {
				if(expy < 0) {
					r = 0x40000000 + ((expy>-30) ? ((r>>-expy)|((r&((static_cast<uint32>(1)<<-expy)-1))!=0)) : 1);
					expy = 0;
				} else {
					r += 0x40000000 >> expy;
					i = r >> 31;
					r = (r>>i) | (r&i);
					expy += i;
				}
			} else {
				r -= 0x40000000 >> expy;
				for(; r<0x40000000; r<<=1,--expy) ;
			}
			my = sqrt<30>(r, expy);
			my = (my<<15) + (r<<14)/my;
			if(S) {
				mx >>= expy - expx;
				ilog = expy;
			} else {
				my >>= expx - expy;
				ilog = expx;
			}
			my += mx;
			i = my >> 31;
			static const int G = S && (R==std::round_to_nearest);
			return log2_post<R,0xB8AA3B2A>(log2(my>>i, 26+S+G)+(G<<3), ilog+i, 17, arg&(static_cast<unsigned>(S)<<15));
		}

		/// Class for 1.31 unsigned floating-point computation
		struct f31 {
			/// Constructor.
			/// \param mant mantissa as 1.31
			/// \param e exponent
			constexpr f31(uint32 mant, int e) : m(mant), exp(e) {}

			/// Constructor.
			/// \param abs unsigned half-precision value
			f31(unsigned int abs) : exp(-15) {
				for(; abs<0x400; abs<<=1,--exp) ;
				m = static_cast<uint32>((abs&0x3FF)|0x400) << 21;
				exp += (abs>>10);
			}

			/// Addition operator.
			/// \param a first operand
			/// \param b second operand
			/// \return \a a + \a b
			friend f31 operator+(f31 a, f31 b) {
				if(b.exp > a.exp)
					std::swap(a, b);
				int d = a.exp - b.exp;
				uint32 m = a.m + ((d<32) ? (b.m>>d) : 0);
				int i = (m&0xFFFFFFFF) < a.m;
				return f31(((m+i)>>i)|0x80000000, a.exp+i);
			}

			/// Subtraction operator.
			/// \param a first operand
			/// \param b second operand
			/// \return \a a - \a b
			friend f31 operator-(f31 a, f31 b) {
				int d = a.exp - b.exp, exp = a.exp;
				uint32 m = a.m - ((d<32) ? (b.m>>d) : 0);
				if(!m)
					return f31(0, -32);
				for(; m<0x80000000; m<<=1,--exp) ;
				return f31(m, exp);
			}

			/// Multiplication operator.
			/// \param a first operand
			/// \param b second operand
			/// \return \a a * \a b
			friend f31 operator*(f31 a, f31 b) {
				uint32 m = multiply64(a.m, b.m);
				int i = m >> 31;
				return f31(m<<(1-i), a.exp + b.exp + i);
			}

			/// Division operator.
			/// \param a first operand
			/// \param b second operand
			/// \return \a a / \a b
			friend f31 operator/(f31 a, f31 b) {
				int i = a.m >= b.m, s;
				uint32 m = divide64((a.m+i)>>i, b.m, s);
				return f31(m, a.exp - b.exp + i - 1);
			}

			uint32 m;			///< mantissa as 1.31.
			int exp;			///< exponent.
		};

		/// Error function and postprocessing.
		/// This computes the value directly in Q1.31 using the approximations given 
		/// [here](https://en.wikipedia.org/wiki/Error_function#Approximation_with_elementary_functions).
		/// \tparam R rounding mode to use
		/// \tparam C `true` for comlementary error function, `false` else
		/// \param arg half-precision function argument
		/// \return approximated value of error function in half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if no other exception occurred
		template<std::float_round_style R,bool C> unsigned int erf(unsigned int arg) {
			unsigned int abs = arg & 0x7FFF, sign = arg & 0x8000;
			f31 x(abs), x2 = x * x * f31(0xB8AA3B29, 0), t = f31(0x80000000, 0) / (f31(0x80000000, 0)+f31(0xA7BA054A, -2)*x), t2 = t * t;
			f31 e = ((f31(0x87DC2213, 0)*t2+f31(0xB5F0E2AE, 0))*t2+f31(0x82790637, -2)-(f31(0xBA00E2B8, 0)*t2+f31(0x91A98E62, -2))*t) * t /
					((x2.exp<0) ? f31(exp2((x2.exp>-32) ? (x2.m>>-x2.exp) : 0, 30), 0) : f31(exp2((x2.m<<x2.exp)&0x7FFFFFFF, 22), x2.m>>(31-x2.exp)));
			return (!C || sign) ? fixed2half<R,31,false,true,true>(0x80000000-(e.m>>(C-e.exp)), 14+C, sign&(C-1U)) :
					(e.exp<-25) ? underflow<R>() : fixed2half<R,30,false,false,true>(e.m>>1, e.exp+14, 0, e.m&1);
		}

		/// Gamma function and postprocessing.
		/// This approximates the value of either the gamma function or its logarithm directly in Q1.31.
		/// \tparam R rounding mode to use
		/// \tparam L `true` for lograithm of gamma function, `false` for gamma function
		/// \param arg half-precision floating-point value
		/// \return lgamma/tgamma(\a arg) in half-precision
		/// \exception FE_OVERFLOW on overflows
		/// \exception FE_UNDERFLOW on underflows
		/// \exception FE_INEXACT if \a arg is not a positive integer
		template<std::float_round_style R,bool L> unsigned int gamma(unsigned int arg) {
/*			static const double p[] ={ 2.50662827563479526904, 225.525584619175212544, -268.295973841304927459, 80.9030806934622512966, -5.00757863970517583837, 0.0114684895434781459556 };
			double t = arg + 4.65, s = p[0];
			for(unsigned int i=0; i<5; ++i)
				s += p[i+1] / (arg+i);
			return std::log(s) + (arg-0.5)*std::log(t) - t;
*/			static const f31 pi(0xC90FDAA2, 1), lbe(0xB8AA3B29, 0);
			unsigned int abs = arg & 0x7FFF, sign = arg & 0x8000;
			bool bsign = sign != 0;
			f31 z(abs), x = sign ? (z+f31(0x80000000, 0)) : z, t = x + f31(0x94CCCCCD, 2), s =
				f31(0xA06C9901, 1) + f31(0xBBE654E2, -7)/(x+f31(0x80000000, 2)) + f31(0xA1CE6098, 6)/(x+f31(0x80000000, 1))
				+ f31(0xE1868CB7, 7)/x - f31(0x8625E279, 8)/(x+f31(0x80000000, 0)) - f31(0xA03E158F, 2)/(x+f31(0xC0000000, 1));
			int i = (s.exp>=2) + (s.exp>=4) + (s.exp>=8) + (s.exp>=16);
			s = f31((static_cast<uint32>(s.exp)<<(31-i))+(log2(s.m>>1, 28)>>i), i) / lbe;
			if(x.exp != -1 || x.m != 0x80000000) {
				i = (t.exp>=2) + (t.exp>=4) + (t.exp>=8);
				f31 l = f31((static_cast<uint32>(t.exp)<<(31-i))+(log2(t.m>>1, 30)>>i), i) / lbe;
				s = (x.exp<-1) ? (s-(f31(0x80000000, -1)-x)*l) : (s+(x-f31(0x80000000, -1))*l);
			}
			s = x.exp ? (s-t) : (t-s);
			if(bsign) {
				if(z.exp >= 0) {
					sign &= (L|((z.m>>(31-z.exp))&1)) - 1;
					for(z=f31((z.m<<(1+z.exp))&0xFFFFFFFF, -1); z.m<0x80000000; z.m<<=1,--z.exp) ;
				}
				if(z.exp == -1)
					z = f31(0x80000000, 0) - z;
				if(z.exp < -1) {
					z = z * pi;
					z.m = sincos(z.m>>(1-z.exp), 30).first;
					for(z.exp=1; z.m<0x80000000; z.m<<=1,--z.exp) ;
				}
				else
					z = f31(0x80000000, 0);
			} if(L) {
				if(bsign) {
					f31 l(0x92868247, 0);
					if(z.exp < 0) {
						uint32 m = log2((z.m+1)>>1, 27);
						z = f31(-((static_cast<uint32>(z.exp)<<26)+(m>>5)), 5);
						for(; z.m<0x80000000; z.m<<=1,--z.exp) ;
						l = l + z / lbe;
					}
					sign = static_cast<unsigned>(x.exp&&(l.exp<s.exp||(l.exp==s.exp&&l.m<s.m))) << 15;
					s = sign ? (s-l) : x.exp ? (l-s) : (l+s);
				} else {
					sign = static_cast<unsigned>(x.exp==0) << 15;
					if(s.exp < -24)
						return underflow<R>(sign);
					if(s.exp > 15)
						return overflow<R>(sign);
				}
			} else {
				s = s * lbe;
				uint32 m;
				if(s.exp < 0) {
					m = s.m >> -s.exp;
					s.exp = 0;
				} else {
					m = (s.m<<s.exp) & 0x7FFFFFFF;
					s.exp = (s.m>>(31-s.exp));
				}
				s.m = exp2(m, 27);
				if(!x.exp)
					s = f31(0x80000000, 0) / s;
				if(bsign) {
					if(z.exp < 0)
						s = s * z;
					s = pi / s;
					if(s.exp < -24)
						return underflow<R>(sign);
				} else if(z.exp > 0 && !(z.m&((1<<(31-z.exp))-1)))
					return ((s.exp+14)<<10) + (s.m>>21);
				if(s.exp > 15)
					return overflow<R>(sign);
			}
			return fixed2half<R,31,false,false,true>(s.m, s.exp+14, sign);
		}
		/// \}

		template<class,class,std::float_round_style> struct half_caster;
	}

	/// Half-precision floating-point type.
	/// This class implements an IEEE-conformant half-precision floating-point type with the usual arithmetic 
	/// operators and conversions. It is implicitly convertible to single-precision floating-point, which makes artihmetic 
	/// expressions and functions with mixed-type operands to be of the most precise operand type.
	///
	/// According to the C++98/03 definition, the half type is not a POD type. But according to C++11's less strict and 
	/// extended definitions it is both a standard layout type and a trivially copyable type (even if not a POD type), which 
	/// means it can be standard-conformantly copied using raw binary copies. But in this context some more words about the 
	/// actual size of the type. Although the half is representing an IEEE 16-bit type, it does not neccessarily have to be of 
	/// exactly 16-bits size. But on any reasonable implementation the actual binary representation of this type will most 
	/// probably not ivolve any additional "magic" or padding beyond the simple binary representation of the underlying 16-bit 
	/// IEEE number, even if not strictly guaranteed by the standard. But even then it only has an actual size of 16 bits if 
	/// your C++ implementation supports an unsigned integer type of exactly 16 bits width. But this should be the case on 
	/// nearly any reasonable platform.
	///
	/// So if your C++ implementation is not totally exotic or imposes special alignment requirements, it is a reasonable 
	/// assumption that the data of a half is just comprised of the 2 bytes of the underlying IEEE representation.
	class half {
	public:
		/// \name Construction and assignment
		/// \{

		/// Default constructor.
		/// This initializes the half to 0. Although this does not match the builtin types' default-initialization semantics 
		/// and may be less efficient than no initialization, it is needed to provide proper value-initialization semantics.
		constexpr half() noexcept : data_() {}

		/// Conversion constructor.
		/// \param rhs float to convert
		/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
		//explicit half(float rhs) : data_(static_cast<detail::uint16>(detail::float2half<round_style>(rhs))) {}

		/// Conversion constructor.
		/// \param rhs float to convert
		/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
		template<class T>
		half(T rhs) : data_(static_cast<detail::uint16>(detail::float2half<round_style>(static_cast<float>(rhs)))) {}

		/// Conversion to single-precision.
		/// \return single precision value representing expression value
		operator float() const { return detail::half2float<float>(data_); }

		/// Assignment operator.
		/// \param rhs single-precision value to copy from
		/// \return reference to this half
		/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
		half& operator=(const float &rhs) { data_ = static_cast<detail::uint16>(detail::float2half<round_style>(rhs)); return *this; }

		template<class T>
		half& operator=(const T &rhs) { return *this = static_cast<float>(rhs); }

		/// \}
		/// \name Arithmetic updates
		/// \{

		/// Arithmetic assignment.
		/// \tparam T type of concrete half expression
		/// \param rhs half expression to add
		/// \return reference to this half
		/// \exception FE_... according to operator+(half,half)
		half& operator+=(half rhs) { return *this = *this + rhs; }

		/// Arithmetic assignment.
		/// \tparam T type of concrete half expression
		/// \param rhs half expression to subtract
		/// \return reference to this half
		/// \exception FE_... according to operator-(half,half)
		half& operator-=(half rhs) { return *this = *this - rhs; }

		/// Arithmetic assignment.
		/// \tparam T type of concrete half expression
		/// \param rhs half expression to multiply with
		/// \return reference to this half
		/// \exception FE_... according to operator*(half,half)
		half& operator*=(half rhs) { return *this = *this * rhs; }

		/// Arithmetic assignment.
		/// \tparam T type of concrete half expression
		/// \param rhs half expression to divide by
		/// \return reference to this half
		/// \exception FE_... according to operator/(half,half)
		half& operator/=(half rhs) { return *this = *this / rhs; }

		/*
		/// Arithmetic assignment.
		/// \param rhs single-precision value to add
		/// \return reference to this half
		/// \exception FE_... according to operator=()
		half& operator+=(float rhs) { return *this = *this + rhs; }

		/// Arithmetic assignment.
		/// \param rhs single-precision value to subtract
		/// \return reference to this half
		/// \exception FE_... according to operator=()
		half& operator-=(float rhs) { return *this = *this - rhs; }

		/// Arithmetic assignment.
		/// \param rhs single-precision value to multiply with
		/// \return reference to this half
		/// \exception FE_... according to operator=()
		half& operator*=(float rhs) { return *this = *this * rhs; }

		/// Arithmetic assignment.
		/// \param rhs single-precision value to divide by
		/// \return reference to this half
		/// \exception FE_... according to operator=()
		half& operator/=(float rhs) { return *this = *this / rhs; }
		*/

		/// \}
		/// \name Increment and decrement
		/// \{

		/// Prefix increment.
		/// \return incremented half value
		/// \exception FE_... according to operator+(half,half)
		half& operator++() { return *this = *this + half(detail::binary, 0x3C00); }

		/// Prefix decrement.
		/// \return decremented half value
		/// \exception FE_... according to operator-(half,half)
		half& operator--() { return *this = *this + half(detail::binary, 0xBC00); }

		/// Postfix increment.
		/// \return non-incremented half value
		/// \exception FE_... according to operator+(half,half)
		half operator++(int) { half out(*this); ++*this; return out; }

		/// Postfix decrement.
		/// \return non-decremented half value
		/// \exception FE_... according to operator-(half,half)
		half operator--(int) { half out(*this); --*this; return out; }
		/// \}
		detail::uint16 get_data()const{ return data_; }
	
	private:
		/// Rounding mode to use
		static const std::float_round_style round_style = (std::float_round_style)(HALF_ROUND_STYLE);

		/// Constructor.
		/// \param bits binary representation to set half to
		constexpr half(detail::binary_t, unsigned int bits) noexcept : data_(static_cast<detail::uint16>(bits)) {}

		/// Internal binary representation
		detail::uint16 data_;

		friend constexpr_NOERR bool operator==(half, half);
		template<class T> friend constexpr_NOERR bool operator==(half, T);
		template<class T> friend constexpr_NOERR bool operator==(T, half);
		friend constexpr_NOERR bool operator!=(half, half);
		template<class T> friend constexpr_NOERR bool operator!=(half, T);
		template<class T> friend constexpr_NOERR bool operator!=(T, half);
		friend constexpr_NOERR bool operator<(half, half);
		template<class T> friend constexpr_NOERR bool operator<(half, T);
		template<class T> friend constexpr_NOERR bool operator<(T, half);
		friend constexpr_NOERR bool operator>(half, half);
		template<class T> friend constexpr_NOERR bool operator>(half, T);
		template<class T> friend constexpr_NOERR bool operator>(T, half);
		friend constexpr_NOERR bool operator<=(half, half);
		template<class T> friend constexpr_NOERR bool operator<=(half, T);
		template<class T> friend constexpr_NOERR bool operator<=(T, half);
		friend constexpr_NOERR bool operator>=(half, half);
		template<class T> friend constexpr_NOERR bool operator>=(half, T);
		template<class T> friend constexpr_NOERR bool operator>=(T, half);
		friend constexpr half operator+(half);
		friend constexpr half operator-(half);
		friend half operator+(half, half);
		template<class T> friend half operator+(half, T);
		template<class T> friend half operator+(T, half);
		friend half operator-(half, half);
		template<class T> friend half operator-(half, T);
		template<class T> friend half operator-(T, half);
		friend half operator*(half, half);
		template<class T> friend half operator*(half, T);
		template<class T> friend half operator*(T, half);
		friend half operator/(half, half);
		template<class T> friend half operator/(half, T);
		template<class T> friend half operator/(T, half);
		template<class charT,class traits> friend std::basic_ostream<charT,traits>& operator<<(std::basic_ostream<charT,traits>&, half);
		template<class charT,class traits> friend std::basic_istream<charT,traits>& operator>>(std::basic_istream<charT,traits>&, half&);
		friend constexpr half fabs(half);
		friend half fmod(half, half);
		friend half remainder(half, half);
		friend half remquo(half, half, int*);
		friend half fma(half, half, half);
		friend constexpr_NOERR half fmax(half, half);
		friend constexpr_NOERR half fmin(half, half);
		friend half fdim(half, half);
		friend half nanh(const char*);
		friend half exp(half);
		friend half exp2(half);
		friend half expm1(half);
		friend half log(half);
		friend half log10(half);
		friend half log2(half);
		friend half log1p(half);
		friend half sqrt(half);
		friend half cbrt(half);
		friend half hypot(half, half);
		friend half hypot(half, half, half);
		friend half pow(half, half);
		friend void sincos(half, half*, half*);
		friend half sin(half);
		friend half cos(half);
		friend half tan(half);
		friend half asin(half);
		friend half acos(half);
		friend half atan(half);
		friend half atan2(half, half);
		friend half sinh(half);
		friend half cosh(half);
		friend half tanh(half);
		friend half asinh(half);
		friend half acosh(half);
		friend half atanh(half);
		friend half erf(half);
		friend half erfc(half);
		friend half lgamma(half);
		friend half tgamma(half);
		friend half ceil(half);
		friend half floor(half);
		friend half trunc(half);
		friend half round(half);
		friend long lround(half);
		friend half rint(half);
		friend long lrint(half);
		friend half nearbyint(half);
		friend long long llround(half);
		friend long long llrint(half);
		friend half frexp(half, int*);
		friend half scalbln(half, long);
		friend half modf(half, half*);
		friend int ilogb(half);
		friend half logb(half);
		friend half nextafter(half, half);
		friend half nexttoward(half, long double);
		friend constexpr half copysign(half, half);
		friend constexpr int fpclassify(half);
		friend constexpr bool isfinite(half);
		friend constexpr bool isinf(half);
		friend constexpr bool isnan(half);
		friend constexpr bool isnormal(half);
		friend constexpr bool signbit(half);
		friend constexpr bool isgreater(half, half);
		friend constexpr bool isgreaterequal(half, half);
		friend constexpr bool isless(half, half);
		friend constexpr bool islessequal(half, half);
		friend constexpr bool islessgreater(half, half);
		template<class,class,std::float_round_style> friend struct detail::half_caster;
		friend class std::numeric_limits<half>;
		friend struct std::hash<half>;
		friend half literal::operator "" _h(long double);
	};

	namespace literal {
		/// Half literal.
		/// While this returns a properly rounded half-precision value, half literals can unfortunately not be constant 
		/// expressions due to rather involved conversions. So don't expect this to be a literal literal without involving 
		/// conversion operations at runtime. It is a convenience feature, not a performance optimization.
		/// \param value literal value
		/// \return half with of given value (possibly rounded)
		/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
		inline half operator "" _h(long double value) { return half(detail::binary, detail::float2half<half::round_style>(value)); }
	}

	namespace detail {
		/// Helper class for half casts.
		/// This class template has to be specialized for all valid cast arguments to define an appropriate static 
		/// `cast` member function and a corresponding `type` member denoting its return type.
		/// \tparam T destination type
		/// \tparam U source type
		/// \tparam R rounding mode to use
		template<class T,class U,std::float_round_style R=(std::float_round_style)(HALF_ROUND_STYLE)> struct half_caster {};
		template<class U,std::float_round_style R> struct half_caster<half,U,R> {
			static_assert(std::is_arithmetic<U>::value, "half_cast from non-arithmetic type unsupported");
			static half cast(U arg) { return cast_impl(arg, is_float<U>()); };
		private:
			static half cast_impl(U arg, true_type) { return half(binary, float2half<R>(arg)); }
			static half cast_impl(U arg, false_type) { return half(binary, int2half<R>(arg)); }
		};
		template<class T,std::float_round_style R> struct half_caster<T,half,R> {
			static_assert(std::is_arithmetic<T>::value, "half_cast to non-arithmetic type unsupported");
			static T cast(half arg) { return cast_impl(arg, is_float<T>()); }
		private:
			static T cast_impl(half arg, true_type) { return half2float<T>(arg.data_); }
			static T cast_impl(half arg, false_type) { return half2int<R,true,true,T>(arg.data_); }
		};
		template<std::float_round_style R> struct half_caster<half,half,R> {
			static half cast(half arg) { return arg; }
		};
	}
}

/// Extensions to the C++ standard library.
namespace std {
	/// Numeric limits for half-precision floats.
	/// **See also:** Documentation for [std::numeric_limits](https://en.cppreference.com/w/cpp/types/numeric_limits)
	template<> class numeric_limits<half_float::half> {
	public:
		/// Is template specialization.
		static constexpr bool is_specialized = true;

		/// Supports signed values.
		static constexpr bool is_signed = true;

		/// Is not an integer type.
		static constexpr bool is_integer = false;

		/// Is not exact.
		static constexpr bool is_exact = false;

		/// Doesn't provide modulo arithmetic.
		static constexpr bool is_modulo = false;

		/// Has a finite set of values.
		static constexpr bool is_bounded = true;

		/// IEEE conformant.
		static constexpr bool is_iec559 = true;

		/// Supports infinity.
		static constexpr bool has_infinity = true;

		/// Supports quiet NaNs.
		static constexpr bool has_quiet_NaN = true;

		/// Supports signaling NaNs.
		static constexpr bool has_signaling_NaN = true;

		/// Supports subnormal values.
		static constexpr float_denorm_style has_denorm = denorm_present;

		/// Supports no denormalization detection.
		static constexpr bool has_denorm_loss = false;

	#if HALF_ERRHANDLING_THROWS
		static constexpr bool traps = true;
	#else
		/// Traps only if [HALF_ERRHANDLING_THROW_...](\ref HALF_ERRHANDLING_THROW_INVALID) is acitvated.
		static constexpr bool traps = false;
	#endif

		/// Does not support no pre-rounding underflow detection.
		static constexpr bool tinyness_before = false;

		/// Rounding mode.
		static constexpr float_round_style round_style = half_float::half::round_style;

		/// Significant digits.
		static constexpr int digits = 11;

		/// Significant decimal digits.
		static constexpr int digits10 = 3;

		/// Required decimal digits to represent all possible values.
		static constexpr int max_digits10 = 5;

		/// Number base.
		static constexpr int radix = 2;

		/// One more than smallest exponent.
		static constexpr int min_exponent = -13;

		/// Smallest normalized representable power of 10.
		static constexpr int min_exponent10 = -4;

		/// One more than largest exponent
		static constexpr int max_exponent = 16;

		/// Largest finitely representable power of 10.
		static constexpr int max_exponent10 = 4;

		/// Smallest positive normal value.
		static constexpr half_float::half min() noexcept { return half_float::half(half_float::detail::binary, 0x0400); }

		/// Smallest finite value.
		static constexpr half_float::half lowest() noexcept { return half_float::half(half_float::detail::binary, 0xFBFF); }

		/// Largest finite value.
		static constexpr half_float::half max() noexcept { return half_float::half(half_float::detail::binary, 0x7BFF); }

		/// Difference between 1 and next representable value.
		static constexpr half_float::half epsilon() noexcept { return half_float::half(half_float::detail::binary, 0x1400); }

		/// Maximum rounding error in ULP (units in the last place).
		static constexpr half_float::half round_error() noexcept
			{ return half_float::half(half_float::detail::binary, (round_style==std::round_to_nearest) ? 0x3800 : 0x3C00); }

		/// Positive infinity.
		static constexpr half_float::half infinity() noexcept { return half_float::half(half_float::detail::binary, 0x7C00); }

		/// Quiet NaN.
		static constexpr half_float::half quiet_NaN() noexcept { return half_float::half(half_float::detail::binary, 0x7FFF); }

		/// Signaling NaN.
		static constexpr half_float::half signaling_NaN() noexcept { return half_float::half(half_float::detail::binary, 0x7DFF); }

		/// Smallest positive subnormal value.
		static constexpr half_float::half denorm_min() noexcept { return half_float::half(half_float::detail::binary, 0x0001); }
	};

	/// Hash function for half-precision floats.
	/// **See also:** Documentation for [std::hash](https://en.cppreference.com/w/cpp/utility/hash)
	template<> struct hash<half_float::half> {
		/// Type of function argument.
		typedef half_float::half argument_type;

		/// Function return type.
		typedef size_t result_type;

		/// Compute hash function.
		/// \param arg half to hash
		/// \return hash value
		result_type operator()(argument_type arg) const { return hash<half_float::detail::uint16>()(arg.data_&-static_cast<unsigned>(arg.data_!=0x8000)); }
	};
}

namespace half_float {
	/// \anchor compop
	/// \name Comparison operators
	/// \{

	/// Comparison for equality.
	/// \param x first operand
	/// \param y second operand
	/// \retval true if operands equal
	/// \retval false else
	/// \exception FE_INVALID if \a x or \a y is NaN
	inline constexpr_NOERR bool operator==(half x, half y) {
		return !detail::compsignal(x.data_, y.data_) && (x.data_==y.data_ || !((x.data_|y.data_)&0x7FFF));
	}
	template<class T>
	inline constexpr_NOERR bool operator==(half x, T y) { return x == static_cast<half>(y); }
	template<class T>
	inline constexpr_NOERR bool operator==(T x, half y) { return static_cast<half>(x) == y; }

	/// Comparison for inequality.
	/// \param x first operand
	/// \param y second operand
	/// \retval true if operands not equal
	/// \retval false else
	/// \exception FE_INVALID if \a x or \a y is NaN
	inline constexpr_NOERR bool operator!=(half x, half y) {
		return detail::compsignal(x.data_, y.data_) || (x.data_!=y.data_ && ((x.data_|y.data_)&0x7FFF));
	}
	template<class T>
	inline constexpr_NOERR bool operator!=(half x, T y) { return x != static_cast<half>(y); }
	template<class T>
	inline constexpr_NOERR bool operator!=(T x, half y) { return static_cast<half>(x) != y; }

	/// Comparison for less than.
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x less than \a y
	/// \retval false else
	/// \exception FE_INVALID if \a x or \a y is NaN
	inline constexpr_NOERR bool operator<(half x, half y) {
		return !detail::compsignal(x.data_, y.data_) &&
			((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) < ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15));
	}
	template<class T>
	inline constexpr_NOERR bool operator<(half x, T y) { return x < static_cast<half>(y); }
	template<class T>
	inline constexpr_NOERR bool operator<(T x, half y) { return static_cast<half>(x) < y; }

	/// Comparison for greater than.
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x greater than \a y
	/// \retval false else
	/// \exception FE_INVALID if \a x or \a y is NaN
	inline constexpr_NOERR bool operator>(half x, half y) {
		return !detail::compsignal(x.data_, y.data_) &&
			((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) > ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15));
	}
	template<class T>
	inline constexpr_NOERR bool operator>(half x, T y) { return x > static_cast<half>(y); }
	template<class T>
	inline constexpr_NOERR bool operator>(T x, half y) { return static_cast<half>(x) > y; }

	/// Comparison for less equal.
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x less equal \a y
	/// \retval false else
	/// \exception FE_INVALID if \a x or \a y is NaN
	inline constexpr_NOERR bool operator<=(half x, half y) {
		return !detail::compsignal(x.data_, y.data_) &&
			((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) <= ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15));
	}
	template<class T>
	inline constexpr_NOERR bool operator<=(half x, T y) { return x <= static_cast<half>(y); }
	template<class T>
	inline constexpr_NOERR bool operator<=(T x, half y) { return static_cast<half>(x) <= y; }

	/// Comparison for greater equal.
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x greater equal \a y
	/// \retval false else
	/// \exception FE_INVALID if \a x or \a y is NaN
	inline constexpr_NOERR bool operator>=(half x, half y) {
		return !detail::compsignal(x.data_, y.data_) &&
			((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) >= ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15));
	}
	template<class T>
	inline constexpr_NOERR bool operator>=(half x, T y) { return x >= static_cast<half>(y); }
	template<class T>
	inline constexpr_NOERR bool operator>=(T x, half y) { return static_cast<half>(x) >= y; }

	/// \}
	/// \anchor arithmetics
	/// \name Arithmetic operators
	/// \{

	/// Identity.
	/// \param arg operand
	/// \return unchanged operand
	inline constexpr half operator+(half arg) { return arg; }

	/// Negation.
	/// \param arg operand
	/// \return negated operand
	inline constexpr half operator-(half arg) { return half(detail::binary, arg.data_^0x8000); }

	/// Addition.
	/// This operation is exact to rounding for all rounding modes.
	/// \param x left operand
	/// \param y right operand
	/// \return sum of half expressions
	/// \exception FE_INVALID if \a x and \a y are infinities with different signs or signaling NaNs
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half operator+(half x, half y) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(detail::half2float<detail::internal_t>(x.data_)+detail::half2float<detail::internal_t>(y.data_)));
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF;
		bool sub = ((x.data_^y.data_)&0x8000) != 0;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) : (absy!=0x7C00) ? x.data_ :
										(sub && absx==0x7C00) ? detail::invalid() : y.data_);
		if(!absx)
			return absy ? y : half(detail::binary, (half::round_style==std::round_toward_neg_infinity) ? (x.data_|y.data_) : (x.data_&y.data_));
		if(!absy)
			return x;
		unsigned int sign = ((sub && absy>absx) ? y.data_ : x.data_) & 0x8000;
		if(absy > absx)
			std::swap(absx, absy);
		int exp = (absx>>10) + (absx<=0x3FF), d = exp - (absy>>10) - (absy<=0x3FF), mx = ((absx&0x3FF)|((absx>0x3FF)<<10)) << 3, my;
		if(d < 13) {
			my = ((absy&0x3FF)|((absy>0x3FF)<<10)) << 3;
			my = (my>>d) | ((my&((1<<d)-1))!=0);
		} else
			my = 1;
		if(sub) {
			if(!(mx-=my))
				return half(detail::binary, static_cast<unsigned>(half::round_style==std::round_toward_neg_infinity)<<15);
			for(; mx<0x2000 && exp>1; mx<<=1,--exp) ;
		} else {
			mx += my;
			int i = mx >> 14;
			if((exp+=i) > 30)
				return half(detail::binary, detail::overflow<half::round_style>(sign));
			mx = (mx>>i) | (mx&i);
		}
		return half(detail::binary, detail::rounded<half::round_style,false>(sign+((exp-1)<<10)+(mx>>3), (mx>>2)&1, (mx&0x3)!=0));
	#endif
	}
	template<class T>
	inline half operator+(half x, T y) { return x + static_cast<half>(y); }
	template<class T>
	inline half operator+(T x, half y) { return static_cast<half>(x) + y; }

	/// Subtraction.
	/// This operation is exact to rounding for all rounding modes.
	/// \param x left operand
	/// \param y right operand
	/// \return difference of half expressions
	/// \exception FE_INVALID if \a x and \a y are infinities with equal signs or signaling NaNs
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half operator-(half x, half y) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(detail::half2float<detail::internal_t>(x.data_)-detail::half2float<detail::internal_t>(y.data_)));
	#else
		return x + (-y);
	#endif
	}
	template<class T>
	inline half operator-(half x, T y) { return x - static_cast<half>(y); }
	template<class T>
	inline half operator-(T x, half y) { return static_cast<half>(x) - y; }

	/// Multiplication.
	/// This operation is exact to rounding for all rounding modes.
	/// \param x left operand
	/// \param y right operand
	/// \return product of half expressions
	/// \exception FE_INVALID if multiplying 0 with infinity or if \a x or \a y is signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half operator*(half x, half y) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(detail::half2float<detail::internal_t>(x.data_)*detail::half2float<detail::internal_t>(y.data_)));
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, exp = -16;
		unsigned int sign = (x.data_^y.data_) & 0x8000;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) :
										((absx==0x7C00 && !absy)||(absy==0x7C00 && !absx)) ? detail::invalid() : (sign|0x7C00));
		if(!absx || !absy)
			return half(detail::binary, sign);
		for(; absx<0x400; absx<<=1,--exp) ;
		for(; absy<0x400; absy<<=1,--exp) ;
		detail::uint32 m = static_cast<detail::uint32>((absx&0x3FF)|0x400) * static_cast<detail::uint32>((absy&0x3FF)|0x400);
		int i = m >> 21, s = m & i;
		exp += (absx>>10) + (absy>>10) + i;
		if(exp > 29)
			return half(detail::binary, detail::overflow<half::round_style>(sign));
		else if(exp < -11)
			return half(detail::binary, detail::underflow<half::round_style>(sign));
		return half(detail::binary, detail::fixed2half<half::round_style,20,false,false,false>(m>>i, exp, sign, s));
	#endif
	}
	template<class T>
	inline half operator*(half x, T y) { return x * static_cast<half>(y); }
	template<class T>
	inline half operator*(T x, half y) { return static_cast<half>(x) * y; }

	/// Division.
	/// This operation is exact to rounding for all rounding modes.
	/// \param x left operand
	/// \param y right operand
	/// \return quotient of half expressions
	/// \exception FE_INVALID if dividing 0s or infinities with each other or if \a x or \a y is signaling NaN
	/// \exception FE_DIVBYZERO if dividing finite value by 0
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half operator/(half x, half y) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(detail::half2float<detail::internal_t>(x.data_)/detail::half2float<detail::internal_t>(y.data_)));
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, exp = 14;
		unsigned int sign = (x.data_^y.data_) & 0x8000;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) :
										(absx==absy) ? detail::invalid() : (sign|((absx==0x7C00) ? 0x7C00 : 0)));
		if(!absx)
			return half(detail::binary, absy ? sign : detail::invalid());
		if(!absy)
			return half(detail::binary, detail::pole(sign));
		for(; absx<0x400; absx<<=1,--exp) ;
		for(; absy<0x400; absy<<=1,++exp) ;
		detail::uint32 mx = (absx&0x3FF) | 0x400, my = (absy&0x3FF) | 0x400;
		int i = mx < my;
		exp += (absx>>10) - (absy>>10) - i;
		if(exp > 29)
			return half(detail::binary, detail::overflow<half::round_style>(sign));
		else if(exp < -11)
			return half(detail::binary, detail::underflow<half::round_style>(sign));
		mx <<= 12 + i;
		my <<= 1;
		return half(detail::binary, detail::fixed2half<half::round_style,11,false,false,false>(mx/my, exp, sign, mx%my!=0));
	#endif
	}
	template<class T>
	inline half operator/(half x, T y) { return x / static_cast<half>(y); }
	template<class T>
	inline half operator/(T x, half y) { return static_cast<half>(x) / y; }

	/// \}
	/// \anchor streaming
	/// \name Input and output
	/// \{

	/// Output operator.
	///	This uses the built-in functionality for streaming out floating-point numbers.
	/// \param out output stream to write into
	/// \param arg half expression to write
	/// \return reference to output stream
	template<class charT,class traits> std::basic_ostream<charT,traits>& operator<<(std::basic_ostream<charT,traits> &out, half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return out << detail::half2float<detail::internal_t>(arg.data_);
	#else
		return out << detail::half2float<float>(arg.data_);
	#endif
	}

	/// Input operator.
	///	This uses the built-in functionality for streaming in floating-point numbers, specifically double precision floating 
	/// point numbers (unless overridden with [HALF_ARITHMETIC_TYPE](\ref HALF_ARITHMETIC_TYPE)). So the input string is first 
	/// rounded to double precision using the underlying platform's current floating-point rounding mode before being rounded 
	/// to half-precision using the library's half-precision rounding mode.
	/// \param in input stream to read from
	/// \param arg half to read into
	/// \return reference to input stream
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	template<class charT,class traits> std::basic_istream<charT,traits>& operator>>(std::basic_istream<charT,traits> &in, half &arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		detail::internal_t f;
	#else
		double f;
	#endif
		if(in >> f)
			arg.data_ = detail::float2half<half::round_style>(f);
		return in;
	}

	/// \}
	/// \anchor basic
	/// \name Basic mathematical operations
	/// \{

	/// Absolute value.
	/// **See also:** Documentation for [std::fabs](https://en.cppreference.com/w/cpp/numeric/math/fabs).
	/// \param arg operand
	/// \return absolute value of \a arg
	inline constexpr half fabs(half arg) { return half(detail::binary, arg.data_&0x7FFF); }

	/// Absolute value.
	/// **See also:** Documentation for [std::abs](https://en.cppreference.com/w/cpp/numeric/math/fabs).
	/// \param arg operand
	/// \return absolute value of \a arg
	inline constexpr half abs(half arg) { return fabs(arg); }

	/// Remainder of division.
	/// **See also:** Documentation for [std::fmod](https://en.cppreference.com/w/cpp/numeric/math/fmod).
	/// \param x first operand
	/// \param y second operand
	/// \return remainder of floating-point division.
	/// \exception FE_INVALID if \a x is infinite or \a y is 0 or if \a x or \a y is signaling NaN
	inline half fmod(half x, half y) {
		unsigned int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, sign = x.data_ & 0x8000;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) :
										(absx==0x7C00) ? detail::invalid() : x.data_);
		if(!absy)
			return half(detail::binary, detail::invalid());
		if(!absx)
			return x;
		if(absx == absy)
			return half(detail::binary, sign);
		return half(detail::binary, sign|detail::mod<false,false>(absx, absy));
	}

	/// Remainder of division.
	/// **See also:** Documentation for [std::remainder](https://en.cppreference.com/w/cpp/numeric/math/remainder).
	/// \param x first operand
	/// \param y second operand
	/// \return remainder of floating-point division.
	/// \exception FE_INVALID if \a x is infinite or \a y is 0 or if \a x or \a y is signaling NaN
	inline half remainder(half x, half y) {
		unsigned int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, sign = x.data_ & 0x8000;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) :
										(absx==0x7C00) ? detail::invalid() : x.data_);
		if(!absy)
			return half(detail::binary, detail::invalid());
		if(absx == absy)
			return half(detail::binary, sign);
		return half(detail::binary, sign^detail::mod<false,true>(absx, absy));
	}

	/// Remainder of division.
	/// **See also:** Documentation for [std::remquo](https://en.cppreference.com/w/cpp/numeric/math/remquo).
	/// \param x first operand
	/// \param y second operand
	/// \param quo address to store some bits of quotient at
	/// \return remainder of floating-point division.
	/// \exception FE_INVALID if \a x is infinite or \a y is 0 or if \a x or \a y is signaling NaN
	inline half remquo(half x, half y, int *quo) {
		unsigned int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, value = x.data_ & 0x8000;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) :
										(absx==0x7C00) ? detail::invalid() : (*quo = 0, x.data_));
		if(!absy)
			return half(detail::binary, detail::invalid());
		bool qsign = ((value^y.data_)&0x8000) != 0;
		int q = 1;
		if(absx != absy)
			value ^= detail::mod<true, true>(absx, absy, &q);
		return *quo = qsign ? -q : q, half(detail::binary, value);
	}

	/// Fused multiply add.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::fma](https://en.cppreference.com/w/cpp/numeric/math/fma).
	/// \param x first operand
	/// \param y second operand
	/// \param z third operand
	/// \return ( \a x * \a y ) + \a z rounded as one operation.
	/// \exception FE_INVALID according to operator*() and operator+() unless any argument is a quiet NaN and no argument is a signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding the final addition
	inline half fma(half x, half y, half z) {
	#ifdef HALF_ARITHMETIC_TYPE
		detail::internal_t fx = detail::half2float<detail::internal_t>(x.data_), fy = detail::half2float<detail::internal_t>(y.data_), fz = detail::half2float<detail::internal_t>(z.data_);
		#if FP_FAST_FMA
			return half(detail::binary, detail::float2half<half::round_style>(std::fma(fx, fy, fz)));
		#else
			return half(detail::binary, detail::float2half<half::round_style>(fx*fy+fz));
		#endif
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, absz = z.data_ & 0x7FFF, exp = -15;
		unsigned int sign = (x.data_^y.data_) & 0x8000;
		bool sub = ((sign^z.data_)&0x8000) != 0;
		if(absx >= 0x7C00 || absy >= 0x7C00 || absz >= 0x7C00)
			return	(absx>0x7C00 || absy>0x7C00 || absz>0x7C00) ? half(detail::binary, detail::signal(x.data_, y.data_, z.data_)) :
					(absx==0x7C00) ? half(detail::binary, (!absy || (sub && absz==0x7C00)) ? detail::invalid() : (sign|0x7C00)) :
					(absy==0x7C00) ? half(detail::binary, (!absx || (sub && absz==0x7C00)) ? detail::invalid() : (sign|0x7C00)) : z;
		if(!absx || !absy)
			return absz ? z : half(detail::binary, (half::round_style==std::round_toward_neg_infinity) ? (z.data_|sign) : (z.data_&sign));
		for(; absx<0x400; absx<<=1,--exp) ;
		for(; absy<0x400; absy<<=1,--exp) ;
		detail::uint32 m = static_cast<detail::uint32>((absx&0x3FF)|0x400) * static_cast<detail::uint32>((absy&0x3FF)|0x400);
		int i = m >> 21;
		exp += (absx>>10) + (absy>>10) + i;
		m <<= 3 - i;
		if(absz) {
			int expz = 0;
			for(; absz<0x400; absz<<=1,--expz) ;
			expz += absz >> 10;
			detail::uint32 mz = static_cast<detail::uint32>((absz&0x3FF)|0x400) << 13;
			if(expz > exp || (expz == exp && mz > m)) {
				std::swap(m, mz);
				std::swap(exp, expz);
				if(sub)
					sign = z.data_ & 0x8000;
			}
			int d = exp - expz;
			mz = (d<23) ? ((mz>>d)|((mz&((static_cast<detail::uint32>(1)<<d)-1))!=0)) : 1;
			if(sub) {
				m = m - mz;
				if(!m)
					return half(detail::binary, static_cast<unsigned>(half::round_style==std::round_toward_neg_infinity)<<15);
				for(; m<0x800000; m<<=1,--exp) ;
			} else {
				m += mz;
				i = m >> 24;
				m = (m>>i) | (m&i);
				exp += i;
			}
		}
		if(exp > 30)
			return half(detail::binary, detail::overflow<half::round_style>(sign));
		else if(exp < -10)
			return half(detail::binary, detail::underflow<half::round_style>(sign));
		return half(detail::binary, detail::fixed2half<half::round_style,23,false,false,false>(m, exp-1, sign));
	#endif
	}

	/// Maximum of half expressions.
	/// **See also:** Documentation for [std::fmax](https://en.cppreference.com/w/cpp/numeric/math/fmax).
	/// \param x first operand
	/// \param y second operand
	/// \return maximum of operands, ignoring quiet NaNs
	/// \exception FE_INVALID if \a x or \a y is signaling NaN
	inline constexpr_NOERR half fmax(half x, half y) {
		return half(detail::binary, (!isnan(y) && (isnan(x) || (x.data_^(0x8000|(0x8000-(x.data_>>15)))) < 
			(y.data_^(0x8000|(0x8000-(y.data_>>15)))))) ? detail::select(y.data_, x.data_) : detail::select(x.data_, y.data_));
	}

	/// Minimum of half expressions.
	/// **See also:** Documentation for [std::fmin](https://en.cppreference.com/w/cpp/numeric/math/fmin).
	/// \param x first operand
	/// \param y second operand
	/// \return minimum of operands, ignoring quiet NaNs
	/// \exception FE_INVALID if \a x or \a y is signaling NaN
	inline constexpr_NOERR half fmin(half x, half y) {
		return half(detail::binary, (!isnan(y) && (isnan(x) || (x.data_^(0x8000|(0x8000-(x.data_>>15)))) >
			(y.data_^(0x8000|(0x8000-(y.data_>>15)))))) ? detail::select(y.data_, x.data_) : detail::select(x.data_, y.data_));
	}

	/// Positive difference.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::fdim](https://en.cppreference.com/w/cpp/numeric/math/fdim).
	/// \param x first operand
	/// \param y second operand
	/// \return \a x - \a y or 0 if difference negative
	/// \exception FE_... according to operator-(half,half)
	inline half fdim(half x, half y) {
		if(isnan(x) || isnan(y))
			return half(detail::binary, detail::signal(x.data_, y.data_));
		return (x.data_^(0x8000|(0x8000-(x.data_>>15)))) <= (y.data_^(0x8000|(0x8000-(y.data_>>15)))) ? half(detail::binary, 0) : (x-y);
	}

	/// Get NaN value.
	/// **See also:** Documentation for [std::nan](https://en.cppreference.com/w/cpp/numeric/math/nan).
	/// \param arg string code
	/// \return quiet NaN
	inline half nanh(const char *arg) {
		unsigned int value = 0x7FFF;
		while(*arg)
			value ^= static_cast<unsigned>(*arg++) & 0xFF;
		return half(detail::binary, value);
	}

	/// \}
	/// \anchor exponential
	/// \name Exponential functions
	/// \{

	/// Exponential function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::exp](https://en.cppreference.com/w/cpp/numeric/math/exp).
	/// \param arg function argument
	/// \return e raised to \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half exp(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::exp(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF;
		if(!abs)
			return half(detail::binary, 0x3C00);
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? (0x7C00&((arg.data_>>15)-1U)) : detail::signal(arg.data_));
		if(abs >= 0x4C80)
			return half(detail::binary, (arg.data_&0x8000) ? detail::underflow<half::round_style>() : detail::overflow<half::round_style>());
		detail::uint32 m = detail::multiply64(static_cast<detail::uint32>((abs&0x3FF)+((abs>0x3FF)<<10))<<21, 0xB8AA3B29);
		int e = (abs>>10) + (abs<=0x3FF), exp;
		if(e < 14) {
			exp = 0;
			m >>= 14 - e;
		} else {
			exp = m >> (45-e);
			m = (m<<(e-14)) & 0x7FFFFFFF;
		}
		return half(detail::binary, detail::exp2_post<half::round_style,true>(detail::exp2(m, 26), exp, (arg.data_&0x8000)!=0));
	#endif
	}

	/// Binary exponential.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::exp2](https://en.cppreference.com/w/cpp/numeric/math/exp2).
	/// \param arg function argument
	/// \return 2 raised to \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half exp2(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::exp2(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF;
		if(!abs)
			return half(detail::binary, 0x3C00);
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? (0x7C00&((arg.data_>>15)-1U)) : detail::signal(arg.data_));
		if(abs >= 0x4E40)
			return half(detail::binary, (arg.data_&0x8000) ? detail::underflow<half::round_style>() : detail::overflow<half::round_style>());
		int e = (abs>>10) + (abs<=0x3FF), exp = (abs&0x3FF) + ((abs>0x3FF)<<10);
		detail::uint32 m = detail::exp2((static_cast<detail::uint32>(exp)<<(6+e))&0x7FFFFFFF, 28);
		exp >>= 25 - e;
		if(m == 0x80000000) {
			if(arg.data_&0x8000)
				exp = -exp;
			else if(exp > 15)
				return half(detail::binary, detail::overflow<half::round_style>());
			return half(detail::binary, detail::fixed2half<half::round_style,31,false,false,false>(m, exp+14));
		}
		return half(detail::binary, detail::exp2_post<half::round_style,true>(m, exp, (arg.data_&0x8000)!=0));
	#endif
	}

	/// Exponential minus one.
	/// This function may be 1 ULP off the correctly rounded exact result in <0.05% of inputs for `std::round_to_nearest` 
	/// and in <1% of inputs for any other rounding mode.
	///
	/// **See also:** Documentation for [std::expm1](https://en.cppreference.com/w/cpp/numeric/math/expm1).
	/// \param arg function argument
	/// \return e raised to \a arg and subtracted by 1
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half expm1(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::expm1(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF, sign = arg.data_ & 0x8000;
		if(!abs)
			return arg;
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? (0x7C00+(sign>>1)) : detail::signal(arg.data_));
		if(abs >= 0x4A00)
			return half(detail::binary, (arg.data_&0x8000) ? detail::rounded<half::round_style,true>(0xBBFF, 1, 1) : detail::overflow<half::round_style>());
		detail::uint32 m = detail::multiply64(static_cast<detail::uint32>((abs&0x3FF)+((abs>0x3FF)<<10))<<21, 0xB8AA3B29);
		int e = (abs>>10) + (abs<=0x3FF), exp;
		if(e < 14) {
			exp = 0;
			m >>= 14 - e;
		} else {
			exp = m >> (45-e);
			m = (m<<(e-14)) & 0x7FFFFFFF;
		}
		m = detail::exp2(m);
		if(sign) {
			int s = 0;
			if(m > 0x80000000) {
				++exp;
				m = detail::divide64(0x80000000, m, s);
			}
			m = 0x80000000 - ((m>>exp)|((m&((static_cast<detail::uint32>(1)<<exp)-1))!=0)|s);
			exp = 0;
		} else
			m -= (exp<31) ? (0x80000000>>exp) : 1;
		for(exp+=14; m<0x80000000 && exp; m<<=1,--exp) ;
		if(exp > 29)
			return half(detail::binary, detail::overflow<half::round_style>());
		return half(detail::binary, detail::rounded<half::round_style,true>(sign+(exp<<10)+(m>>21), (m>>20)&1, (m&0xFFFFF)!=0));
	#endif
	}

	/// Natural logarithm.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::log](https://en.cppreference.com/w/cpp/numeric/math/log).
	/// \param arg function argument
	/// \return logarithm of \a arg to base e
	/// \exception FE_INVALID for signaling NaN or negative argument
	/// \exception FE_DIVBYZERO for 0
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half log(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::log(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = -15;
		if(!abs)
			return half(detail::binary, detail::pole(0x8000));
		if(arg.data_ & 0x8000)
			return half(detail::binary, (arg.data_<=0xFC00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs >= 0x7C00)
			return (abs==0x7C00) ? arg : half(detail::binary, detail::signal(arg.data_));
		for(; abs<0x400; abs<<=1,--exp) ;
		exp += abs >> 10;
		return half(detail::binary, detail::log2_post<half::round_style,0xB8AA3B2A>(
			detail::log2(static_cast<detail::uint32>((abs&0x3FF)|0x400)<<20, 27)+8, exp, 17));
	#endif
	}

	/// Common logarithm.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::log10](https://en.cppreference.com/w/cpp/numeric/math/log10).
	/// \param arg function argument
	/// \return logarithm of \a arg to base 10
	/// \exception FE_INVALID for signaling NaN or negative argument
	/// \exception FE_DIVBYZERO for 0
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half log10(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::log10(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = -15;
		if(!abs)
			return half(detail::binary, detail::pole(0x8000));
		if(arg.data_ & 0x8000)
			return half(detail::binary, (arg.data_<=0xFC00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs >= 0x7C00)
			return (abs==0x7C00) ? arg : half(detail::binary, detail::signal(arg.data_));
		switch(abs) {
			case 0x4900: return half(detail::binary, 0x3C00);
			case 0x5640: return half(detail::binary, 0x4000);
			case 0x63D0: return half(detail::binary, 0x4200);
			case 0x70E2: return half(detail::binary, 0x4400);
		}
		for(; abs<0x400; abs<<=1,--exp) ;
		exp += abs >> 10;
		return half(detail::binary, detail::log2_post<half::round_style,0xD49A784C>(
			detail::log2(static_cast<detail::uint32>((abs&0x3FF)|0x400)<<20, 27)+8, exp, 16));
	#endif
	}

	/// Binary logarithm.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::log2](https://en.cppreference.com/w/cpp/numeric/math/log2).
	/// \param arg function argument
	/// \return logarithm of \a arg to base 2
	/// \exception FE_INVALID for signaling NaN or negative argument
	/// \exception FE_DIVBYZERO for 0
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half log2(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::log2(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = -15, s = 0;
		if(!abs)
			return half(detail::binary, detail::pole(0x8000));
		if(arg.data_ & 0x8000)
			return half(detail::binary, (arg.data_<=0xFC00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs >= 0x7C00)
			return (abs==0x7C00) ? arg : half(detail::binary, detail::signal(arg.data_));
		if(abs == 0x3C00)
			return half(detail::binary, 0);
		for(; abs<0x400; abs<<=1,--exp) ;
		exp += (abs>>10);
		if(!(abs&0x3FF)) {
			unsigned int value = static_cast<unsigned>(exp<0) << 15, m = std::abs(exp) << 6;
			for(exp=18; m<0x400; m<<=1,--exp) ;
			return half(detail::binary, value+(exp<<10)+m);
		}
		detail::uint32 ilog = exp, sign = detail::sign_mask(ilog), m = 
			(((ilog<<27)+(detail::log2(static_cast<detail::uint32>((abs&0x3FF)|0x400)<<20, 28)>>4))^sign) - sign;
		if(!m)
			return half(detail::binary, 0);
		for(exp=14; m<0x8000000 && exp; m<<=1,--exp) ;
		for(; m>0xFFFFFFF; m>>=1,++exp)
			s |= m & 1;
		return half(detail::binary, detail::fixed2half<half::round_style,27,false,false,true>(m, exp, sign&0x8000, s));
	#endif
	}

	/// Natural logarithm plus one.
	/// This function may be 1 ULP off the correctly rounded exact result in <0.05% of inputs for `std::round_to_nearest` 
	/// and in ~1% of inputs for any other rounding mode.
	///
	/// **See also:** Documentation for [std::log1p](https://en.cppreference.com/w/cpp/numeric/math/log1p).
	/// \param arg function argument
	/// \return logarithm of \a arg plus 1 to base e
	/// \exception FE_INVALID for signaling NaN or argument <-1
	/// \exception FE_DIVBYZERO for -1
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half log1p(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::log1p(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		if(arg.data_ >= 0xBC00)
			return half(detail::binary, (arg.data_==0xBC00) ? detail::pole(0x8000) : (arg.data_<=0xFC00) ? detail::invalid() : detail::signal(arg.data_));
		int abs = arg.data_ & 0x7FFF, exp = -15;
		if(!abs || abs >= 0x7C00)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		for(; abs<0x400; abs<<=1,--exp) ;
		exp += abs >> 10;
		detail::uint32 m = static_cast<detail::uint32>((abs&0x3FF)|0x400) << 20;
		if(arg.data_ & 0x8000) {
			m = 0x40000000 - (m>>-exp);
			for(exp=0; m<0x40000000; m<<=1,--exp) ;
		} else {
			if(exp < 0) {
				m = 0x40000000 + (m>>-exp);
				exp = 0;
			} else {
				m += 0x40000000 >> exp;
				int i = m >> 31;
				m >>= i;
				exp += i;
			}
		}
		return half(detail::binary, detail::log2_post<half::round_style,0xB8AA3B2A>(detail::log2(m), exp, 17));
	#endif
	}

	/// \}
	/// \anchor power
	/// \name Power functions
	/// \{

	/// Square root.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::sqrt](https://en.cppreference.com/w/cpp/numeric/math/sqrt).
	/// \param arg function argument
	/// \return square root of \a arg
	/// \exception FE_INVALID for signaling NaN and negative arguments
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half sqrt(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::sqrt(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = 15;
		if(!abs || arg.data_ >= 0x7C00)
			return half(detail::binary, (abs>0x7C00) ? detail::signal(arg.data_) : (arg.data_>0x8000) ? detail::invalid() : arg.data_);
		for(; abs<0x400; abs<<=1,--exp) ;
		detail::uint32 r = static_cast<detail::uint32>((abs&0x3FF)|0x400) << 10, m = detail::sqrt<20>(r, exp+=abs>>10);
		return half(detail::binary, detail::rounded<half::round_style,false>((exp<<10)+(m&0x3FF), r>m, r!=0));
	#endif
	}

	/// Cubic root.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::cbrt](https://en.cppreference.com/w/cpp/numeric/math/cbrt).
	/// \param arg function argument
	/// \return cubic root of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half cbrt(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::cbrt(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = -15;
		if(!abs || abs == 0x3C00 || abs >= 0x7C00)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		for(; abs<0x400; abs<<=1, --exp);
		detail::uint32 ilog = exp + (abs>>10), sign = detail::sign_mask(ilog), f, m = 
			(((ilog<<27)+(detail::log2(static_cast<detail::uint32>((abs&0x3FF)|0x400)<<20, 24)>>4))^sign) - sign;
		for(exp=2; m<0x80000000; m<<=1,--exp) ;
		m = detail::multiply64(m, 0xAAAAAAAB);
		int i = m >> 31, s;
		exp += i;
		m <<= 1 - i;
		if(exp < 0) {
			f = m >> -exp;
			exp = 0;
		} else {
			f = (m<<exp) & 0x7FFFFFFF;
			exp = m >> (31-exp);
		}
		m = detail::exp2(f, (half::round_style==std::round_to_nearest) ? 29 : 26);
		if(sign) {
			if(m > 0x80000000) {
				m = detail::divide64(0x80000000, m, s);
				++exp;
			}
			exp = -exp;
		}
		return half(detail::binary, (half::round_style==std::round_to_nearest) ?
			detail::fixed2half<half::round_style,31,false,false,false>(m, exp+14, arg.data_&0x8000) :
			detail::fixed2half<half::round_style,23,false,false,false>((m+0x80)>>8, exp+14, arg.data_&0x8000));
	#endif
	}

	/// Hypotenuse function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::hypot](https://en.cppreference.com/w/cpp/numeric/math/hypot).
	/// \param x first argument
	/// \param y second argument
	/// \return square root of sum of squares without internal over- or underflows
	/// \exception FE_INVALID if \a x or \a y is signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding of the final square root
	inline half hypot(half x, half y) {
	#ifdef HALF_ARITHMETIC_TYPE
		detail::internal_t fx = detail::half2float<detail::internal_t>(x.data_), fy = detail::half2float<detail::internal_t>(y.data_);
		return half(detail::binary, detail::float2half<half::round_style>(std::hypot(fx, fy)));
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, expx = 0, expy = 0;
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx==0x7C00) ? detail::select(0x7C00, y.data_) :
				(absy==0x7C00) ? detail::select(0x7C00, x.data_) : detail::signal(x.data_, y.data_));
		if(!absx)
			return half(detail::binary, absy ? detail::check_underflow(absy) : 0);
		if(!absy)
			return half(detail::binary, detail::check_underflow(absx));
		if(absy > absx)
			std::swap(absx, absy);
		for(; absx<0x400; absx<<=1,--expx) ;
		for(; absy<0x400; absy<<=1,--expy) ;
		detail::uint32 mx = (absx&0x3FF) | 0x400, my = (absy&0x3FF) | 0x400;
		mx *= mx;
		my *= my;
		int ix = mx >> 21, iy = my >> 21;
		expx = 2*(expx+(absx>>10)) - 15 + ix;
		expy = 2*(expy+(absy>>10)) - 15 + iy;
		mx <<= 10 - ix;
		my <<= 10 - iy;
		int d = expx - expy;
		my = (d<30) ? ((my>>d)|((my&((static_cast<detail::uint32>(1)<<d)-1))!=0)) : 1;
		return half(detail::binary, detail::hypot_post<half::round_style>(mx+my, expx));
	#endif
	}

	/// Hypotenuse function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::hypot](https://en.cppreference.com/w/cpp/numeric/math/hypot).
	/// \param x first argument
	/// \param y second argument
	/// \param z third argument
	/// \return square root of sum of squares without internal over- or underflows
	/// \exception FE_INVALID if \a x, \a y or \a z is signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding of the final square root
	inline half hypot(half x, half y, half z) {
	#ifdef HALF_ARITHMETIC_TYPE
		detail::internal_t fx = detail::half2float<detail::internal_t>(x.data_), fy = detail::half2float<detail::internal_t>(y.data_), fz = detail::half2float<detail::internal_t>(z.data_);
		return half(detail::binary, detail::float2half<half::round_style>(std::sqrt(fx*fx+fy*fy+fz*fz)));
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, absz = z.data_ & 0x7FFF, expx = 0, expy = 0, expz = 0;
		if(!absx)
			return hypot(y, z);
		if(!absy)
			return hypot(x, z);
		if(!absz)
			return hypot(x, y);
		if(absx >= 0x7C00 || absy >= 0x7C00 || absz >= 0x7C00)
			return half(detail::binary,	(absx==0x7C00) ? detail::select(0x7C00, detail::select(y.data_, z.data_)) :
										(absy==0x7C00) ? detail::select(0x7C00, detail::select(x.data_, z.data_)) :
										(absz==0x7C00) ? detail::select(0x7C00, detail::select(x.data_, y.data_)) :
										detail::signal(x.data_, y.data_, z.data_));
		if(absz > absy)
			std::swap(absy, absz);
		if(absy > absx)
			std::swap(absx, absy);
		if(absz > absy)
			std::swap(absy, absz);
		for(; absx<0x400; absx<<=1,--expx) ;
		for(; absy<0x400; absy<<=1,--expy) ;
		for(; absz<0x400; absz<<=1,--expz) ;
		detail::uint32 mx = (absx&0x3FF) | 0x400, my = (absy&0x3FF) | 0x400, mz = (absz&0x3FF) | 0x400;
		mx *= mx;
		my *= my;
		mz *= mz;
		int ix = mx >> 21, iy = my >> 21, iz = mz >> 21;
		expx = 2*(expx+(absx>>10)) - 15 + ix;
		expy = 2*(expy+(absy>>10)) - 15 + iy;
		expz = 2*(expz+(absz>>10)) - 15 + iz;
		mx <<= 10 - ix;
		my <<= 10 - iy;
		mz <<= 10 - iz;
		int d = expy - expz;
		mz = (d<30) ? ((mz>>d)|((mz&((static_cast<detail::uint32>(1)<<d)-1))!=0)) : 1;
		my += mz;
		if(my & 0x80000000) {
			my = (my>>1) | (my&1);
			if(++expy > expx) {
				std::swap(mx, my);
				std::swap(expx, expy);
			}
		}
		d = expx - expy;
		my = (d<30) ? ((my>>d)|((my&((static_cast<detail::uint32>(1)<<d)-1))!=0)) : 1;
		return half(detail::binary, detail::hypot_post<half::round_style>(mx+my, expx));
	#endif
	}

	/// Power function.
	/// This function may be 1 ULP off the correctly rounded exact result for any rounding mode in ~0.00025% of inputs.
	///
	/// **See also:** Documentation for [std::pow](https://en.cppreference.com/w/cpp/numeric/math/pow).
	/// \param x base
	/// \param y exponent
	/// \return \a x raised to \a y
	/// \exception FE_INVALID if \a x or \a y is signaling NaN or if \a x is finite an negative and \a y is finite and not integral
	/// \exception FE_DIVBYZERO if \a x is 0 and \a y is negative
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half pow(half x, half y) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::pow(detail::half2float<detail::internal_t>(x.data_), detail::half2float<detail::internal_t>(y.data_))));
	#else
		int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, exp = -15;
		if(!absy || x.data_ == 0x3C00)
			return half(detail::binary, detail::select(0x3C00, (x.data_==0x3C00) ? y.data_ : x.data_));
		bool is_int = absy >= 0x6400 || (absy>=0x3C00 && !(absy&((1<<(25-(absy>>10)))-1)));
		unsigned int sign = x.data_ & (static_cast<unsigned>((absy<0x6800)&&is_int&&((absy>>(25-(absy>>10)))&1))<<15);
		if(absx >= 0x7C00 || absy >= 0x7C00)
			return half(detail::binary,	(absx>0x7C00 || absy>0x7C00) ? detail::signal(x.data_, y.data_) :
										(absy==0x7C00) ? ((absx==0x3C00) ? 0x3C00 : (!absx && y.data_==0xFC00) ? detail::pole() :
										(0x7C00&-((y.data_>>15)^(absx>0x3C00)))) : (sign|(0x7C00&((y.data_>>15)-1U))));
		if(!absx)
			return half(detail::binary, (y.data_&0x8000) ? detail::pole(sign) : sign);
		if((x.data_&0x8000) && !is_int)
			return half(detail::binary, detail::invalid());
		if(x.data_ == 0xBC00)
			return half(detail::binary, sign|0x3C00);
		if(y.data_ == 0x3800)
			return sqrt(x);
		if(y.data_ == 0x3C00)
			return half(detail::binary, detail::check_underflow(x.data_));
		if(y.data_ == 0x4000)
			return x * x;
		for(; absx<0x400; absx<<=1,--exp) ;
		detail::uint32 ilog = exp + (absx>>10), msign = detail::sign_mask(ilog), f, m = 
			(((ilog<<27)+((detail::log2(static_cast<detail::uint32>((absx&0x3FF)|0x400)<<20)+8)>>4))^msign) - msign;
		for(exp=-11; m<0x80000000; m<<=1,--exp) ;
		for(; absy<0x400; absy<<=1,--exp) ;
		m = detail::multiply64(m, static_cast<detail::uint32>((absy&0x3FF)|0x400)<<21);
		int i = m >> 31;
		exp += (absy>>10) + i;
		m <<= 1 - i;
		if(exp < 0) {
			f = m >> -exp;
			exp = 0;
		} else {
			f = (m<<exp) & 0x7FFFFFFF;
			exp = m >> (31-exp);
		}
		return half(detail::binary, detail::exp2_post<half::round_style,false>(detail::exp2(f), exp, ((msign&1)^(y.data_>>15))!=0, sign));
	#endif
	}

	/// \}
	/// \anchor trigonometric
	/// \name Trigonometric functions
	/// \{

	/// Compute sine and cosine simultaneously.
	///	This returns the same results as sin() and cos() but is faster than calling each function individually.
	///
	/// This function is exact to rounding for all rounding modes.
	/// \param arg function argument
	/// \param sin variable to take sine of \a arg
	/// \param cos variable to take cosine of \a arg
	/// \exception FE_INVALID for signaling NaN or infinity
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline void sincos(half arg, half *sin, half *cos) {
	#ifdef HALF_ARITHMETIC_TYPE
		detail::internal_t f = detail::half2float<detail::internal_t>(arg.data_);
		*sin = half(detail::binary, detail::float2half<half::round_style>(std::sin(f)));
		*cos = half(detail::binary, detail::float2half<half::round_style>(std::cos(f)));
	#else
		int abs = arg.data_ & 0x7FFF, sign = arg.data_ >> 15, k;
		if(abs >= 0x7C00)
			*sin = *cos = half(detail::binary, (abs==0x7C00) ? detail::invalid() : detail::signal(arg.data_));
		else if(!abs) {
			*sin = arg;
			*cos = half(detail::binary, 0x3C00);
		} else if(abs < 0x2500) {
			*sin = half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-1, 1, 1));
			*cos = half(detail::binary, detail::rounded<half::round_style,true>(0x3BFF, 1, 1));
		} else {
			if(half::round_style != std::round_to_nearest) {
				switch(abs) {
				case 0x48B7:
					*sin = half(detail::binary, detail::rounded<half::round_style,true>((~arg.data_&0x8000)|0x1D07, 1, 1));
					*cos = half(detail::binary, detail::rounded<half::round_style,true>(0xBBFF, 1, 1));
					return;
				case 0x598C:
					*sin = half(detail::binary, detail::rounded<half::round_style,true>((arg.data_&0x8000)|0x3BFF, 1, 1));
					*cos = half(detail::binary, detail::rounded<half::round_style,true>(0x80FC, 1, 1));
					return;
				case 0x6A64:
					*sin = half(detail::binary, detail::rounded<half::round_style,true>((~arg.data_&0x8000)|0x3BFE, 1, 1));
					*cos = half(detail::binary, detail::rounded<half::round_style,true>(0x27FF, 1, 1));
					return;
				case 0x6D8C:
					*sin = half(detail::binary, detail::rounded<half::round_style,true>((arg.data_&0x8000)|0x0FE6, 1, 1));
					*cos = half(detail::binary, detail::rounded<half::round_style,true>(0x3BFF, 1, 1));
					return;
				}
			}
			std::pair<detail::uint32,detail::uint32> sc = detail::sincos(detail::angle_arg(abs, k), 28);
			switch(k & 3) {
				case 1: sc = std::make_pair(sc.second, -sc.first); break;
				case 2: sc = std::make_pair(-sc.first, -sc.second); break;
				case 3: sc = std::make_pair(-sc.second, sc.first); break;
			}
			*sin = half(detail::binary, detail::fixed2half<half::round_style,30,true,true,true>((sc.first^-static_cast<detail::uint32>(sign))+sign));
			*cos = half(detail::binary, detail::fixed2half<half::round_style,30,true,true,true>(sc.second));
		}
	#endif
	}

	/// Sine function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::sin](https://en.cppreference.com/w/cpp/numeric/math/sin).
	/// \param arg function argument
	/// \return sine value of \a arg
	/// \exception FE_INVALID for signaling NaN or infinity
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half sin(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::sin(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, k;
		if(!abs)
			return arg;
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs < 0x2900)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-1, 1, 1));
		if(half::round_style != std::round_to_nearest)
			switch(abs) {
				case 0x48B7: return half(detail::binary, detail::rounded<half::round_style,true>((~arg.data_&0x8000)|0x1D07, 1, 1));
				case 0x6A64: return half(detail::binary, detail::rounded<half::round_style,true>((~arg.data_&0x8000)|0x3BFE, 1, 1));
				case 0x6D8C: return half(detail::binary, detail::rounded<half::round_style,true>((arg.data_&0x8000)|0x0FE6, 1, 1));
			}
		std::pair<detail::uint32,detail::uint32> sc = detail::sincos(detail::angle_arg(abs, k), 28);
		detail::uint32 sign = -static_cast<detail::uint32>(((k>>1)&1)^(arg.data_>>15));
		return half(detail::binary, detail::fixed2half<half::round_style,30,true,true,true>((((k&1) ? sc.second : sc.first)^sign) - sign));
	#endif
	}

	/// Cosine function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::cos](https://en.cppreference.com/w/cpp/numeric/math/cos).
	/// \param arg function argument
	/// \return cosine value of \a arg
	/// \exception FE_INVALID for signaling NaN or infinity
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half cos(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::cos(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, k;
		if(!abs)
			return half(detail::binary, 0x3C00);
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs < 0x2500)
			return half(detail::binary, detail::rounded<half::round_style,true>(0x3BFF, 1, 1));
		if(half::round_style != std::round_to_nearest && abs == 0x598C)
			return half(detail::binary, detail::rounded<half::round_style,true>(0x80FC, 1, 1));
		std::pair<detail::uint32,detail::uint32> sc = detail::sincos(detail::angle_arg(abs, k), 28);
		detail::uint32 sign = -static_cast<detail::uint32>(((k>>1)^k)&1);
		return half(detail::binary, detail::fixed2half<half::round_style,30,true,true,true>((((k&1) ? sc.first : sc.second)^sign) - sign));
	#endif
	}

	/// Tangent function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::tan](https://en.cppreference.com/w/cpp/numeric/math/tan).
	/// \param arg function argument
	/// \return tangent value of \a arg
	/// \exception FE_INVALID for signaling NaN or infinity
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half tan(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::tan(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = 13, k;
		if(!abs)
			return arg;
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs < 0x2700)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_, 0, 1));
		if(half::round_style != std::round_to_nearest)
			switch(abs) {
				case 0x658C: return half(detail::binary, detail::rounded<half::round_style,true>((arg.data_&0x8000)|0x07E6, 1, 1));
				case 0x7330: return half(detail::binary, detail::rounded<half::round_style,true>((~arg.data_&0x8000)|0x4B62, 1, 1));
			}
		std::pair<detail::uint32,detail::uint32> sc = detail::sincos(detail::angle_arg(abs, k), 30);
		if(k & 1)
			sc = std::make_pair(-sc.second, sc.first);
		detail::uint32 signy = detail::sign_mask(sc.first), signx = detail::sign_mask(sc.second);
		detail::uint32 my = (sc.first^signy) - signy, mx = (sc.second^signx) - signx;
		for(; my<0x80000000; my<<=1,--exp) ;
		for(; mx<0x80000000; mx<<=1,++exp) ;
		return half(detail::binary, detail::tangent_post<half::round_style>(my, mx, exp, (signy^signx^arg.data_)&0x8000));
	#endif
	}

	/// Arc sine.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::asin](https://en.cppreference.com/w/cpp/numeric/math/asin).
	/// \param arg function argument
	/// \return arc sine value of \a arg
	/// \exception FE_INVALID for signaling NaN or if abs(\a arg) > 1
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half asin(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::asin(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF, sign = arg.data_ & 0x8000;
		if(!abs)
			return arg;
		if(abs >= 0x3C00)
			return half(detail::binary, (abs>0x7C00) ? detail::signal(arg.data_) : (abs>0x3C00) ? detail::invalid() :
										detail::rounded<half::round_style,true>(sign|0x3E48, 0, 1));
		if(abs < 0x2900)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_, 0, 1));
		if(half::round_style != std::round_to_nearest && (abs == 0x2B44 || abs == 0x2DC3))
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_+1, 1, 1));
		std::pair<detail::uint32,detail::uint32> sc = detail::atan2_args(abs);
		detail::uint32 m = detail::atan2(sc.first, sc.second, (half::round_style==std::round_to_nearest) ? 27 : 26);
		return half(detail::binary, detail::fixed2half<half::round_style,30,false,true,true>(m, 14, sign));
	#endif
	}

	/// Arc cosine function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::acos](https://en.cppreference.com/w/cpp/numeric/math/acos).
	/// \param arg function argument
	/// \return arc cosine value of \a arg
	/// \exception FE_INVALID for signaling NaN or if abs(\a arg) > 1
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half acos(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::acos(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF, sign = arg.data_ >> 15;
		if(!abs)
			return half(detail::binary, detail::rounded<half::round_style,true>(0x3E48, 0, 1));
		if(abs >= 0x3C00)
			return half(detail::binary,	(abs>0x7C00) ? detail::signal(arg.data_) : (abs>0x3C00) ? detail::invalid() :
										sign ? detail::rounded<half::round_style,true>(0x4248, 0, 1) : 0);
		std::pair<detail::uint32,detail::uint32> cs = detail::atan2_args(abs);
		detail::uint32 m = detail::atan2(cs.second, cs.first, 28);
		return half(detail::binary, detail::fixed2half<half::round_style,31,false,true,true>(sign ? (0xC90FDAA2-m) : m, 15, 0, sign));
	#endif
	}

	/// Arc tangent function.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::atan](https://en.cppreference.com/w/cpp/numeric/math/atan).
	/// \param arg function argument
	/// \return arc tangent value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half atan(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::atan(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF, sign = arg.data_ & 0x8000;
		if(!abs)
			return arg;
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? detail::rounded<half::round_style,true>(sign|0x3E48, 0, 1) : detail::signal(arg.data_));
		if(abs <= 0x2700)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-1, 1, 1));
		int exp = (abs>>10) + (abs<=0x3FF);
		detail::uint32 my = (abs&0x3FF) | ((abs>0x3FF)<<10);
		detail::uint32 m = (exp>15) ?	detail::atan2(my<<19, 0x20000000>>(exp-15), (half::round_style==std::round_to_nearest) ? 26 : 24) :
										detail::atan2(my<<(exp+4), 0x20000000, (half::round_style==std::round_to_nearest) ? 30 : 28);
		return half(detail::binary, detail::fixed2half<half::round_style,30,false,true,true>(m, 14, sign));
	#endif
	}

	/// Arc tangent function.
	/// This function may be 1 ULP off the correctly rounded exact result in ~0.005% of inputs for `std::round_to_nearest`, 
	/// in ~0.1% of inputs for `std::round_toward_zero` and in ~0.02% of inputs for any other rounding mode.
	///
	/// **See also:** Documentation for [std::atan2](https://en.cppreference.com/w/cpp/numeric/math/atan2).
	/// \param y numerator
	/// \param x denominator
	/// \return arc tangent value
	/// \exception FE_INVALID if \a x or \a y is signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half atan2(half y, half x) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::atan2(detail::half2float<detail::internal_t>(y.data_), detail::half2float<detail::internal_t>(x.data_))));
	#else
		unsigned int absx = x.data_ & 0x7FFF, absy = y.data_ & 0x7FFF, signx = x.data_ >> 15, signy = y.data_ & 0x8000;
		if(absx >= 0x7C00 || absy >= 0x7C00) {
			if(absx > 0x7C00 || absy > 0x7C00)
				return half(detail::binary, detail::signal(x.data_, y.data_));
			if(absy == 0x7C00)
				return half(detail::binary, (absx<0x7C00) ?	detail::rounded<half::round_style,true>(signy|0x3E48, 0, 1) :
													signx ?	detail::rounded<half::round_style,true>(signy|0x40B6, 0, 1) :
															detail::rounded<half::round_style,true>(signy|0x3A48, 0, 1));
			return (x.data_==0x7C00) ? half(detail::binary, signy) : half(detail::binary, detail::rounded<half::round_style,true>(signy|0x4248, 0, 1));
		}
		if(!absy)
			return signx ? half(detail::binary, detail::rounded<half::round_style,true>(signy|0x4248, 0, 1)) : y;
		if(!absx)
			return half(detail::binary, detail::rounded<half::round_style,true>(signy|0x3E48, 0, 1));
		int d = (absy>>10) + (absy<=0x3FF) - (absx>>10) - (absx<=0x3FF);
		if(d > (signx ? 18 : 12))
			return half(detail::binary, detail::rounded<half::round_style,true>(signy|0x3E48, 0, 1));
		if(signx && d < -11)
			return half(detail::binary, detail::rounded<half::round_style,true>(signy|0x4248, 0, 1));
		if(!signx && d < ((half::round_style==std::round_toward_zero) ? -15 : -9)) {
			for(; absy<0x400; absy<<=1,--d) ;
			detail::uint32 mx = ((absx<<1)&0x7FF) | 0x800, my = ((absy<<1)&0x7FF) | 0x800;
			int i = my < mx;
			d -= i;
			if(d < -25)
				return half(detail::binary, detail::underflow<half::round_style>(signy));
			my <<= 11 + i;
			return half(detail::binary, detail::fixed2half<half::round_style,11,false,false,true>(my/mx, d+14, signy, my%mx!=0));
		}
		detail::uint32 m = detail::atan2(	((absy&0x3FF)|((absy>0x3FF)<<10))<<(19+((d<0) ? d : (d>0) ? 0 : -1)),
											((absx&0x3FF)|((absx>0x3FF)<<10))<<(19-((d>0) ? d : (d<0) ? 0 : 1)));
		return half(detail::binary, detail::fixed2half<half::round_style,31,false,true,true>(signx ? (0xC90FDAA2-m) : m, 15, signy, signx));
	#endif
	}

	/// \}
	/// \anchor hyperbolic
	/// \name Hyperbolic functions
	/// \{

	/// Hyperbolic sine.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::sinh](https://en.cppreference.com/w/cpp/numeric/math/sinh).
	/// \param arg function argument
	/// \return hyperbolic sine value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half sinh(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::sinh(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp;
		if(!abs || abs >= 0x7C00)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		if(abs <= 0x2900)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_, 0, 1));
		std::pair<detail::uint32,detail::uint32> mm = detail::hyperbolic_args(abs, exp, (half::round_style==std::round_to_nearest) ? 29 : 27);
		detail::uint32 m = mm.first - mm.second;
		for(exp+=13; m<0x80000000 && exp; m<<=1,--exp) ;
		unsigned int sign = arg.data_ & 0x8000;
		if(exp > 29)
			return half(detail::binary, detail::overflow<half::round_style>(sign));
		return half(detail::binary, detail::fixed2half<half::round_style,31,false,false,true>(m, exp, sign));
	#endif
	}

	/// Hyperbolic cosine.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::cosh](https://en.cppreference.com/w/cpp/numeric/math/cosh).
	/// \param arg function argument
	/// \return hyperbolic cosine value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half cosh(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::cosh(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp;
		if(!abs)
			return half(detail::binary, 0x3C00);
		if(abs >= 0x7C00)
			return half(detail::binary, (abs>0x7C00) ? detail::signal(arg.data_) : 0x7C00);
		std::pair<detail::uint32,detail::uint32> mm = detail::hyperbolic_args(abs, exp, (half::round_style==std::round_to_nearest) ? 23 : 26);
		detail::uint32 m = mm.first + mm.second, i = (~m&0xFFFFFFFF) >> 31;
		m = (m>>i) | (m&i) | 0x80000000;
		if((exp+=13+i) > 29)
			return half(detail::binary, detail::overflow<half::round_style>());
		return half(detail::binary, detail::fixed2half<half::round_style,31,false,false,true>(m, exp));
	#endif
	}

	/// Hyperbolic tangent.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::tanh](https://en.cppreference.com/w/cpp/numeric/math/tanh).
	/// \param arg function argument
	/// \return hyperbolic tangent value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half tanh(half arg) {
	#ifdef HALF_ARITHMETIC_TYPE
		return half(detail::binary, detail::float2half<half::round_style>(std::tanh(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp;
		if(!abs)
			return arg;
		if(abs >= 0x7C00)
			return half(detail::binary, (abs>0x7C00) ? detail::signal(arg.data_) : (arg.data_-0x4000));
		if(abs >= 0x4500)
			return half(detail::binary, detail::rounded<half::round_style,true>((arg.data_&0x8000)|0x3BFF, 1, 1));
		if(abs < 0x2700)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-1, 1, 1));
		if(half::round_style != std::round_to_nearest && abs == 0x2D3F)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-3, 0, 1));
		std::pair<detail::uint32,detail::uint32> mm = detail::hyperbolic_args(abs, exp, 27);
		detail::uint32 my = mm.first - mm.second - (half::round_style!=std::round_to_nearest), mx = mm.first + mm.second, i = (~mx&0xFFFFFFFF) >> 31;
		for(exp=13; my<0x80000000; my<<=1,--exp) ;
		mx = (mx>>i) | 0x80000000;
		return half(detail::binary, detail::tangent_post<half::round_style>(my, mx, exp-i, arg.data_&0x8000));
	#endif
	}

	/// Hyperbolic area sine.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::asinh](https://en.cppreference.com/w/cpp/numeric/math/asinh).
	/// \param arg function argument
	/// \return area sine value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half asinh(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::asinh(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF;
		if(!abs || abs >= 0x7C00)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		if(abs <= 0x2900)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-1, 1, 1));
		if(half::round_style != std::round_to_nearest)
			switch(abs)
			{
				case 0x32D4: return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-13, 1, 1));
				case 0x3B5B: return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_-197, 1, 1));
			}
		return half(detail::binary, detail::area<half::round_style,true>(arg.data_));
	#endif
	}

	/// Hyperbolic area cosine.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::acosh](https://en.cppreference.com/w/cpp/numeric/math/acosh).
	/// \param arg function argument
	/// \return area cosine value of \a arg
	/// \exception FE_INVALID for signaling NaN or arguments <1
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half acosh(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::acosh(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF;
		if((arg.data_&0x8000) || abs < 0x3C00)
			return half(detail::binary, (abs<=0x7C00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs == 0x3C00)
			return half(detail::binary, 0);
		if(arg.data_ >= 0x7C00)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		return half(detail::binary, detail::area<half::round_style,false>(arg.data_));
	#endif
	}

	/// Hyperbolic area tangent.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::atanh](https://en.cppreference.com/w/cpp/numeric/math/atanh).
	/// \param arg function argument
	/// \return area tangent value of \a arg
	/// \exception FE_INVALID for signaling NaN or if abs(\a arg) > 1
	/// \exception FE_DIVBYZERO for +/-1
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half atanh(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::atanh(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF, exp = 0;
		if(!abs)
			return arg;
		if(abs >= 0x3C00)
			return half(detail::binary, (abs==0x3C00) ? detail::pole(arg.data_&0x8000) : (abs<=0x7C00) ? detail::invalid() : detail::signal(arg.data_));
		if(abs < 0x2700)
			return half(detail::binary, detail::rounded<half::round_style,true>(arg.data_, 0, 1));
		detail::uint32 m = static_cast<detail::uint32>((abs&0x3FF)|((abs>0x3FF)<<10)) << ((abs>>10)+(abs<=0x3FF)+6), my = 0x80000000 + m, mx = 0x80000000 - m;
		for(; mx<0x80000000; mx<<=1,++exp) ;
		int i = my >= mx, s;
		return half(detail::binary, detail::log2_post<half::round_style,0xB8AA3B2A>(detail::log2(
			(detail::divide64(my>>i, mx, s)+1)>>1, 27)+0x10, exp+i-1, 16, arg.data_&0x8000));
	#endif
	}

	/// \}
	/// \anchor special
	/// \name Error and gamma functions
	/// \{

	/// Error function.
	/// This function may be 1 ULP off the correctly rounded exact result for any rounding mode in <0.5% of inputs.
	///
	/// **See also:** Documentation for [std::erf](https://en.cppreference.com/w/cpp/numeric/math/erf).
	/// \param arg function argument
	/// \return error function value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half erf(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::erf(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF;
		if(!abs || abs >= 0x7C00)
			return (abs>=0x7C00) ? half(detail::binary, (abs==0x7C00) ? (arg.data_-0x4000) : detail::signal(arg.data_)) : arg;
		if(abs >= 0x4200)
			return half(detail::binary, detail::rounded<half::round_style,true>((arg.data_&0x8000)|0x3BFF, 1, 1));
		return half(detail::binary, detail::erf<half::round_style,false>(arg.data_));
	#endif
	}

	/// Complementary error function.
	/// This function may be 1 ULP off the correctly rounded exact result for any rounding mode in <0.5% of inputs.
	///
	/// **See also:** Documentation for [std::erfc](https://en.cppreference.com/w/cpp/numeric/math/erfc).
	/// \param arg function argument
	/// \return 1 minus error function value of \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half erfc(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::erfc(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF, sign = arg.data_ & 0x8000;
		if(abs >= 0x7C00)
			return (abs>=0x7C00) ? half(detail::binary, (abs==0x7C00) ? (sign>>1) : detail::signal(arg.data_)) : arg;
		if(!abs)
			return half(detail::binary, 0x3C00);
		if(abs >= 0x4400)
			return half(detail::binary, detail::rounded<half::round_style,true>((sign>>1)-(sign>>15), sign>>15, 1));
		return half(detail::binary, detail::erf<half::round_style,true>(arg.data_));
	#endif
	}

	/// Natural logarithm of gamma function.
	/// This function may be 1 ULP off the correctly rounded exact result for any rounding mode in ~0.025% of inputs.
	///
	/// **See also:** Documentation for [std::lgamma](https://en.cppreference.com/w/cpp/numeric/math/lgamma).
	/// \param arg function argument
	/// \return natural logarith of gamma function for \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_DIVBYZERO for 0 or negative integer arguments
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half lgamma(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::lgamma(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		int abs = arg.data_ & 0x7FFF;
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? 0x7C00 : detail::signal(arg.data_));
		if(!abs || arg.data_ >= 0xE400 || (arg.data_ >= 0xBC00 && !(abs&((1<<(25-(abs>>10)))-1))))
			return half(detail::binary, detail::pole());
		if(arg.data_ == 0x3C00 || arg.data_ == 0x4000)
			return half(detail::binary, 0);
		return half(detail::binary, detail::gamma<half::round_style,true>(arg.data_));
	#endif
	}

	/// Gamma function.
	/// This function may be 1 ULP off the correctly rounded exact result for any rounding mode in <0.25% of inputs.
	///
	/// **See also:** Documentation for [std::tgamma](https://en.cppreference.com/w/cpp/numeric/math/tgamma).
	/// \param arg function argument
	/// \return gamma function value of \a arg
	/// \exception FE_INVALID for signaling NaN, negative infinity or negative integer arguments
	/// \exception FE_DIVBYZERO for 0
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half tgamma(half arg) {
	#if defined(HALF_ARITHMETIC_TYPE)
		return half(detail::binary, detail::float2half<half::round_style>(std::tgamma(detail::half2float<detail::internal_t>(arg.data_))));
	#else
		unsigned int abs = arg.data_ & 0x7FFF;
		if(!abs)
			return half(detail::binary, detail::pole(arg.data_));
		if(abs >= 0x7C00)
			return (arg.data_==0x7C00) ? arg : half(detail::binary, detail::signal(arg.data_));
		if(arg.data_ >= 0xE400 || (arg.data_ >= 0xBC00 && !(abs&((1<<(25-(abs>>10)))-1))))
			return half(detail::binary, detail::invalid());
		if(arg.data_ >= 0xCA80)
			return half(detail::binary, detail::underflow<half::round_style>((1-((abs>>(25-(abs>>10)))&1))<<15));
		if(arg.data_ <= 0x100 || (arg.data_ >= 0x4900 && arg.data_ < 0x8000))
			return half(detail::binary, detail::overflow<half::round_style>());
		if(arg.data_ == 0x3C00)
			return arg;
		return half(detail::binary, detail::gamma<half::round_style,false>(arg.data_));
	#endif
	}

	/// \}
	/// \anchor rounding
	/// \name Rounding
	/// \{

	/// Nearest integer not less than half value.
	/// **See also:** Documentation for [std::ceil](https://en.cppreference.com/w/cpp/numeric/math/ceil).
	/// \param arg half to round
	/// \return nearest integer not less than \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_INEXACT if value had to be rounded
	inline half ceil(half arg) { return half(detail::binary, detail::integral<std::round_toward_infinity,true,true>(arg.data_)); }

	/// Nearest integer not greater than half value.
	/// **See also:** Documentation for [std::floor](https://en.cppreference.com/w/cpp/numeric/math/floor).
	/// \param arg half to round
	/// \return nearest integer not greater than \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_INEXACT if value had to be rounded
	inline half floor(half arg) { return half(detail::binary, detail::integral<std::round_toward_neg_infinity,true,true>(arg.data_)); }

	/// Nearest integer not greater in magnitude than half value.
	/// **See also:** Documentation for [std::trunc](https://en.cppreference.com/w/cpp/numeric/math/trunc).
	/// \param arg half to round
	/// \return nearest integer not greater in magnitude than \a arg
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_INEXACT if value had to be rounded
	inline half trunc(half arg) { return half(detail::binary, detail::integral<std::round_toward_zero,true,true>(arg.data_)); }

	/// Nearest integer.
	/// **See also:** Documentation for [std::round](https://en.cppreference.com/w/cpp/numeric/math/round).
	/// \param arg half to round
	/// \return nearest integer, rounded away from zero in half-way cases
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_INEXACT if value had to be rounded
	inline half round(half arg) { return half(detail::binary, detail::integral<std::round_to_nearest,false,true>(arg.data_)); }

	/// Nearest integer.
	/// **See also:** Documentation for [std::lround](https://en.cppreference.com/w/cpp/numeric/math/round).
	/// \param arg half to round
	/// \return nearest integer, rounded away from zero in half-way cases
	/// \exception FE_INVALID if value is not representable as `long`
	inline long lround(half arg) { return detail::half2int<std::round_to_nearest,false,false,long>(arg.data_); }

	/// Nearest integer using half's internal rounding mode.
	/// **See also:** Documentation for [std::rint](https://en.cppreference.com/w/cpp/numeric/math/rint).
	/// \param arg half expression to round
	/// \return nearest integer using default rounding mode
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_INEXACT if value had to be rounded
	inline half rint(half arg) { return half(detail::binary, detail::integral<half::round_style,true,true>(arg.data_)); }

	/// Nearest integer using half's internal rounding mode.
	/// **See also:** Documentation for [std::lrint](https://en.cppreference.com/w/cpp/numeric/math/rint).
	/// \param arg half expression to round
	/// \return nearest integer using default rounding mode
	/// \exception FE_INVALID if value is not representable as `long`
	/// \exception FE_INEXACT if value had to be rounded
	inline long lrint(half arg) { return detail::half2int<half::round_style,true,true,long>(arg.data_); }

	/// Nearest integer using half's internal rounding mode.
	/// **See also:** Documentation for [std::nearbyint](https://en.cppreference.com/w/cpp/numeric/math/nearbyint).
	/// \param arg half expression to round
	/// \return nearest integer using default rounding mode
	/// \exception FE_INVALID for signaling NaN
	inline half nearbyint(half arg) { return half(detail::binary, detail::integral<half::round_style,true,false>(arg.data_)); }
	/// Nearest integer.
	/// **See also:** Documentation for [std::llround](https://en.cppreference.com/w/cpp/numeric/math/round).
	/// \param arg half to round
	/// \return nearest integer, rounded away from zero in half-way cases
	/// \exception FE_INVALID if value is not representable as `long long`
	inline long long llround(half arg) { return detail::half2int<std::round_to_nearest,false,false,long long>(arg.data_); }

	/// Nearest integer using half's internal rounding mode.
	/// **See also:** Documentation for [std::llrint](https://en.cppreference.com/w/cpp/numeric/math/rint).
	/// \param arg half expression to round
	/// \return nearest integer using default rounding mode
	/// \exception FE_INVALID if value is not representable as `long long`
	/// \exception FE_INEXACT if value had to be rounded
	inline long long llrint(half arg) { return detail::half2int<half::round_style,true,true,long long>(arg.data_); }

	/// \}
	/// \anchor float
	/// \name Floating point manipulation
	/// \{

	/// Decompress floating-point number.
	/// **See also:** Documentation for [std::frexp](https://en.cppreference.com/w/cpp/numeric/math/frexp).
	/// \param arg number to decompress
	/// \param exp address to store exponent at
	/// \return significant in range [0.5, 1)
	/// \exception FE_INVALID for signaling NaN
	inline half frexp(half arg, int *exp) {
		*exp = 0;
		unsigned int abs = arg.data_ & 0x7FFF;
		if(abs >= 0x7C00 || !abs)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		for(; abs<0x400; abs<<=1,--*exp) ;
		*exp += (abs>>10) - 14;
		return half(detail::binary, (arg.data_&0x8000)|0x3800|(abs&0x3FF));
	}

	/// Multiply by power of two.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::scalbln](https://en.cppreference.com/w/cpp/numeric/math/scalbn).
	/// \param arg number to modify
	/// \param exp power of two to multiply with
	/// \return \a arg multplied by 2 raised to \a exp
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half scalbln(half arg, long exp) {
		unsigned int abs = arg.data_ & 0x7FFF, sign = arg.data_ & 0x8000;
		if(abs >= 0x7C00 || !abs)
			return (abs>0x7C00) ? half(detail::binary, detail::signal(arg.data_)) : arg;
		for(; abs<0x400; abs<<=1,--exp) ;
		exp += abs >> 10;
		if(exp > 30)
			return half(detail::binary, detail::overflow<half::round_style>(sign));
		else if(exp < -10)
			return half(detail::binary, detail::underflow<half::round_style>(sign));
		else if(exp > 0)
			return half(detail::binary, sign|(exp<<10)|(abs&0x3FF));
		unsigned int m = (abs&0x3FF) | 0x400;
		return half(detail::binary, detail::rounded<half::round_style,false>(sign|(m>>(1-exp)), (m>>-exp)&1, (m&((1<<-exp)-1))!=0));
	}

	/// Multiply by power of two.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::scalbn](https://en.cppreference.com/w/cpp/numeric/math/scalbn).
	/// \param arg number to modify
	/// \param exp power of two to multiply with
	/// \return \a arg multplied by 2 raised to \a exp
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half scalbn(half arg, int exp) { return scalbln(arg, exp); }

	/// Multiply by power of two.
	/// This function is exact to rounding for all rounding modes.
	///
	/// **See also:** Documentation for [std::ldexp](https://en.cppreference.com/w/cpp/numeric/math/ldexp).
	/// \param arg number to modify
	/// \param exp power of two to multiply with
	/// \return \a arg multplied by 2 raised to \a exp
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	inline half ldexp(half arg, int exp) { return scalbln(arg, exp); }

	/// Extract integer and fractional parts.
	/// **See also:** Documentation for [std::modf](https://en.cppreference.com/w/cpp/numeric/math/modf).
	/// \param arg number to decompress
	/// \param iptr address to store integer part at
	/// \return fractional part
	/// \exception FE_INVALID for signaling NaN
	inline half modf(half arg, half *iptr) {
		unsigned int abs = arg.data_ & 0x7FFF;
		if(abs > 0x7C00) {
			arg = half(detail::binary, detail::signal(arg.data_));
			return *iptr = arg, arg;
		}
		if(abs >= 0x6400)
			return *iptr = arg, half(detail::binary, arg.data_&0x8000);
		if(abs < 0x3C00)
			return iptr->data_ = arg.data_ & 0x8000, arg;
		unsigned int exp = abs >> 10, mask = (1<<(25-exp)) - 1, m = arg.data_ & mask;
		iptr->data_ = arg.data_ & ~mask;
		if(!m)
			return half(detail::binary, arg.data_&0x8000);
		for(; m<0x400; m<<=1,--exp) ;
		return half(detail::binary, (arg.data_&0x8000)|(exp<<10)|(m&0x3FF));
	}

	/// Extract exponent.
	/// **See also:** Documentation for [std::ilogb](https://en.cppreference.com/w/cpp/numeric/math/ilogb).
	/// \param arg number to query
	/// \return floating-point exponent
	/// \retval FP_ILOGB0 for zero
	/// \retval FP_ILOGBNAN for NaN
	/// \retval INT_MAX for infinity
	/// \exception FE_INVALID for 0 or infinite values
	inline int ilogb(half arg) {
		int abs = arg.data_ & 0x7FFF, exp;
		if(!abs || abs >= 0x7C00) {
			detail::raise(FE_INVALID);
			return !abs ? FP_ILOGB0 : (abs==0x7C00) ? INT_MAX : FP_ILOGBNAN;
		}
		for(exp=(abs>>10)-15; abs<0x200; abs<<=1,--exp) ;
		return exp;
	}

	/// Extract exponent.
	/// **See also:** Documentation for [std::logb](https://en.cppreference.com/w/cpp/numeric/math/logb).
	/// \param arg number to query
	/// \return floating-point exponent
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_DIVBYZERO for 0
	inline half logb(half arg) {
		int abs = arg.data_ & 0x7FFF, exp;
		if(!abs)
			return half(detail::binary, detail::pole(0x8000));
		if(abs >= 0x7C00)
			return half(detail::binary, (abs==0x7C00) ? 0x7C00 : detail::signal(arg.data_));
		for(exp=(abs>>10)-15; abs<0x200; abs<<=1,--exp) ;
		unsigned int value = static_cast<unsigned>(exp<0) << 15;
		if(exp) {
			unsigned int m = std::abs(exp) << 6;
			for(exp=18; m<0x400; m<<=1,--exp) ;
			value |= (exp<<10) + m;
		}
		return half(detail::binary, value);
	}

	/// Next representable value.
	/// **See also:** Documentation for [std::nextafter](https://en.cppreference.com/w/cpp/numeric/math/nextafter).
	/// \param from value to compute next representable value for
	/// \param to direction towards which to compute next value
	/// \return next representable value after \a from in direction towards \a to
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW for infinite result from finite argument
	/// \exception FE_UNDERFLOW for subnormal result
	inline half nextafter(half from, half to) {
		int fabs = from.data_ & 0x7FFF, tabs = to.data_ & 0x7FFF;
		if(fabs > 0x7C00 || tabs > 0x7C00)
			return half(detail::binary, detail::signal(from.data_, to.data_));
		if(from.data_ == to.data_ || !(fabs|tabs))
			return to;
		if(!fabs) {
			detail::raise(FE_UNDERFLOW, !HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT);
			return half(detail::binary, (to.data_&0x8000)+1);
		}
		unsigned int out = from.data_ + (((from.data_>>15)^static_cast<unsigned>(
			(from.data_^(0x8000|(0x8000-(from.data_>>15))))<(to.data_^(0x8000|(0x8000-(to.data_>>15))))))<<1) - 1;
		detail::raise(FE_OVERFLOW, fabs<0x7C00 && (out&0x7C00)==0x7C00);
		detail::raise(FE_UNDERFLOW, !HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT && (out&0x7C00)<0x400);
		return half(detail::binary, out);
	}

	/// Next representable value.
	/// **See also:** Documentation for [std::nexttoward](https://en.cppreference.com/w/cpp/numeric/math/nexttoward).
	/// \param from value to compute next representable value for
	/// \param to direction towards which to compute next value
	/// \return next representable value after \a from in direction towards \a to
	/// \exception FE_INVALID for signaling NaN
	/// \exception FE_OVERFLOW for infinite result from finite argument
	/// \exception FE_UNDERFLOW for subnormal result
	inline half nexttoward(half from, long double to) {
		int fabs = from.data_ & 0x7FFF;
		if(fabs > 0x7C00)
			return half(detail::binary, detail::signal(from.data_));
		long double lfrom = static_cast<long double>(from);
		if(detail::builtin_isnan(to) || lfrom == to)
			return half(static_cast<float>(to));
		if(!fabs) {
			detail::raise(FE_UNDERFLOW, !HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT);
			return half(detail::binary, (static_cast<unsigned>(detail::builtin_signbit(to))<<15)+1);
		}
		unsigned int out = from.data_ + (((from.data_>>15)^static_cast<unsigned>(lfrom<to))<<1) - 1;
		detail::raise(FE_OVERFLOW, (out&0x7FFF)==0x7C00);
		detail::raise(FE_UNDERFLOW, !HALF_ERRHANDLING_UNDERFLOW_TO_INEXACT && (out&0x7FFF)<0x400);
		return half(detail::binary, out);
	}

	/// Take sign.
	/// **See also:** Documentation for [std::copysign](https://en.cppreference.com/w/cpp/numeric/math/copysign).
	/// \param x value to change sign for
	/// \param y value to take sign from
	/// \return value equal to \a x in magnitude and to \a y in sign
	inline constexpr half copysign(half x, half y) { return half(detail::binary, x.data_^((x.data_^y.data_)&0x8000)); }

	/// \}
	/// \anchor classification
	/// \name Floating point classification
	/// \{

	/// Classify floating-point value.
	/// **See also:** Documentation for [std::fpclassify](https://en.cppreference.com/w/cpp/numeric/math/fpclassify).
	/// \param arg number to classify
	/// \retval FP_ZERO for positive and negative zero
	/// \retval FP_SUBNORMAL for subnormal numbers
	/// \retval FP_INFINITY for positive and negative infinity
	/// \retval FP_NAN for NaNs
	/// \retval FP_NORMAL for all other (normal) values
	inline constexpr int fpclassify(half arg) {
		return	!(arg.data_&0x7FFF) ? FP_ZERO :
				((arg.data_&0x7FFF)<0x400) ? FP_SUBNORMAL :
				((arg.data_&0x7FFF)<0x7C00) ? FP_NORMAL :
				((arg.data_&0x7FFF)==0x7C00) ? FP_INFINITE :
				FP_NAN;
	}

	/// Check if finite number.
	/// **See also:** Documentation for [std::isfinite](https://en.cppreference.com/w/cpp/numeric/math/isfinite).
	/// \param arg number to check
	/// \retval true if neither infinity nor NaN
	/// \retval false else
	inline constexpr bool isfinite(half arg) { return (arg.data_&0x7C00) != 0x7C00; }

	/// Check for infinity.
	/// **See also:** Documentation for [std::isinf](https://en.cppreference.com/w/cpp/numeric/math/isinf).
	/// \param arg number to check
	/// \retval true for positive or negative infinity
	/// \retval false else
	inline constexpr bool isinf(half arg) { return (arg.data_&0x7FFF) == 0x7C00; }

	/// Check for NaN.
	/// **See also:** Documentation for [std::isnan](https://en.cppreference.com/w/cpp/numeric/math/isnan).
	/// \param arg number to check
	/// \retval true for NaNs
	/// \retval false else
	inline constexpr bool isnan(half arg) { return (arg.data_&0x7FFF) > 0x7C00; }

	/// Check if normal number.
	/// **See also:** Documentation for [std::isnormal](https://en.cppreference.com/w/cpp/numeric/math/isnormal).
	/// \param arg number to check
	/// \retval true if normal number
	/// \retval false if either subnormal, zero, infinity or NaN
	inline constexpr bool isnormal(half arg) { return ((arg.data_&0x7C00)!=0) & ((arg.data_&0x7C00)!=0x7C00); }

	/// Check sign.
	/// **See also:** Documentation for [std::signbit](https://en.cppreference.com/w/cpp/numeric/math/signbit).
	/// \param arg number to check
	/// \retval true for negative number
	/// \retval false for positive number
	inline constexpr bool signbit(half arg) { return (arg.data_&0x8000) != 0; }

	/// \}
	/// \anchor compfunc
	/// \name Comparison
	/// \{

	/// Quiet comparison for greater than.
	/// **See also:** Documentation for [std::isgreater](https://en.cppreference.com/w/cpp/numeric/math/isgreater).
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x greater than \a y
	/// \retval false else
	inline constexpr bool isgreater(half x, half y) {
		return ((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) > ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15)) && !isnan(x) && !isnan(y);
	}

	/// Quiet comparison for greater equal.
	/// **See also:** Documentation for [std::isgreaterequal](https://en.cppreference.com/w/cpp/numeric/math/isgreaterequal).
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x greater equal \a y
	/// \retval false else
	inline constexpr bool isgreaterequal(half x, half y) {
		return ((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) >= ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15)) && !isnan(x) && !isnan(y);
	}

	/// Quiet comparison for less than.
	/// **See also:** Documentation for [std::isless](https://en.cppreference.com/w/cpp/numeric/math/isless).
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x less than \a y
	/// \retval false else
	inline constexpr bool isless(half x, half y) {
		return ((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) < ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15)) && !isnan(x) && !isnan(y);
	}

	/// Quiet comparison for less equal.
	/// **See also:** Documentation for [std::islessequal](https://en.cppreference.com/w/cpp/numeric/math/islessequal).
	/// \param x first operand
	/// \param y second operand
	/// \retval true if \a x less equal \a y
	/// \retval false else
	inline constexpr bool islessequal(half x, half y) {
		return ((x.data_^(0x8000|(0x8000-(x.data_>>15))))+(x.data_>>15)) <= ((y.data_^(0x8000|(0x8000-(y.data_>>15))))+(y.data_>>15)) && !isnan(x) && !isnan(y);
	}

	/// Quiet comarison for less or greater.
	/// **See also:** Documentation for [std::islessgreater](https://en.cppreference.com/w/cpp/numeric/math/islessgreater).
	/// \param x first operand
	/// \param y second operand
	/// \retval true if either less or greater
	/// \retval false else
	inline constexpr bool islessgreater(half x, half y) {
		return x.data_!=y.data_ && ((x.data_|y.data_)&0x7FFF) && !isnan(x) && !isnan(y);
	}

	/// Quiet check if unordered.
	/// **See also:** Documentation for [std::isunordered](https://en.cppreference.com/w/cpp/numeric/math/isunordered).
	/// \param x first operand
	/// \param y second operand
	/// \retval true if unordered (one or two NaN operands)
	/// \retval false else
	inline constexpr bool isunordered(half x, half y) { return isnan(x) || isnan(y); }

	/// \}
	/// \anchor casting
	/// \name Casting
	/// \{

	/// Cast to or from half-precision floating-point number.
	/// This casts between [half](\ref half_float::half) and any built-in arithmetic type. The values are converted 
	/// directly using the default rounding mode, without any roundtrip over `float` that a `static_cast` would otherwise do.
	///
	/// Using this cast with neither of the two types being a [half](\ref half_float::half) or with any of the two types 
	/// not being a built-in arithmetic type (apart from [half](\ref half_float::half), of course) results in a compiler 
	/// error and casting between [half](\ref half_float::half)s returns the argument unmodified.
	/// \tparam T destination type (half or built-in arithmetic type)
	/// \tparam U source type (half or built-in arithmetic type)
	/// \param arg value to cast
	/// \return \a arg converted to destination type
	/// \exception FE_INVALID if \a T is integer type and result is not representable as \a T
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	template<class T,class U> T half_cast(U arg) { return detail::half_caster<T,U>::cast(arg); }

	/// Cast to or from half-precision floating-point number.
	/// This casts between [half](\ref half_float::half) and any built-in arithmetic type. The values are converted 
	/// directly using the specified rounding mode, without any roundtrip over `float` that a `static_cast` would otherwise do.
	///
	/// Using this cast with neither of the two types being a [half](\ref half_float::half) or with any of the two types 
	/// not being a built-in arithmetic type (apart from [half](\ref half_float::half), of course) results in a compiler 
	/// error and casting between [half](\ref half_float::half)s returns the argument unmodified.
	/// \tparam T destination type (half or built-in arithmetic type)
	/// \tparam R rounding mode to use.
	/// \tparam U source type (half or built-in arithmetic type)
	/// \param arg value to cast
	/// \return \a arg converted to destination type
	/// \exception FE_INVALID if \a T is integer type and result is not representable as \a T
	/// \exception FE_OVERFLOW, ...UNDERFLOW, ...INEXACT according to rounding
	template<class T,std::float_round_style R,class U> T half_cast(U arg) { return detail::half_caster<T,U,R>::cast(arg); }
	/// \}

	/// \}
	/// \anchor errors
	/// \name Error handling
	/// \{

	/// Clear exception flags.
	/// This function works even if [automatic exception flag handling](\ref HALF_ERRHANDLING_FLAGS) is disabled, 
	/// but in that case manual flag management is the only way to raise flags.
	///
	/// **See also:** Documentation for [std::feclearexcept](https://en.cppreference.com/w/cpp/numeric/fenv/feclearexcept).
	/// \param excepts OR of exceptions to clear
	/// \retval 0 all selected flags cleared successfully
	inline int feclearexcept(int excepts) { detail::errflags() &= ~excepts; return 0; }

	/// Test exception flags.
	/// This function works even if [automatic exception flag handling](\ref HALF_ERRHANDLING_FLAGS) is disabled, 
	/// but in that case manual flag management is the only way to raise flags.
	///
	/// **See also:** Documentation for [std::fetestexcept](https://en.cppreference.com/w/cpp/numeric/fenv/fetestexcept).
	/// \param excepts OR of exceptions to test
	/// \return OR of selected exceptions if raised
	inline int fetestexcept(int excepts) { return detail::errflags() & excepts; }

	/// Raise exception flags.
	/// This raises the specified floating point exceptions and also invokes any additional automatic exception handling as 
	/// configured with the [HALF_ERRHANDLIG_...](\ref HALF_ERRHANDLING_ERRNO) preprocessor symbols.
	/// This function works even if [automatic exception flag handling](\ref HALF_ERRHANDLING_FLAGS) is disabled, 
	/// but in that case manual flag management is the only way to raise flags.
	///
	/// **See also:** Documentation for [std::feraiseexcept](https://en.cppreference.com/w/cpp/numeric/fenv/feraiseexcept).
	/// \param excepts OR of exceptions to raise
	/// \retval 0 all selected exceptions raised successfully
	inline int feraiseexcept(int excepts) { detail::errflags() |= excepts; detail::raise(excepts); return 0; }

	/// Save exception flags.
	/// This function works even if [automatic exception flag handling](\ref HALF_ERRHANDLING_FLAGS) is disabled, 
	/// but in that case manual flag management is the only way to raise flags.
	///
	/// **See also:** Documentation for [std::fegetexceptflag](https://en.cppreference.com/w/cpp/numeric/fenv/feexceptflag).
	/// \param flagp adress to store flag state at
	/// \param excepts OR of flags to save
	/// \retval 0 for success
	inline int fegetexceptflag(int *flagp, int excepts) { *flagp = detail::errflags() & excepts; return 0; }

	/// Restore exception flags.
	/// This only copies the specified exception state (including unset flags) without incurring any additional exception handling.
	/// This function works even if [automatic exception flag handling](\ref HALF_ERRHANDLING_FLAGS) is disabled, 
	/// but in that case manual flag management is the only way to raise flags.
	///
	/// **See also:** Documentation for [std::fesetexceptflag](https://en.cppreference.com/w/cpp/numeric/fenv/feexceptflag).
	/// \param flagp adress to take flag state from
	/// \param excepts OR of flags to restore
	/// \retval 0 for success
	inline int fesetexceptflag(const int *flagp, int excepts) { detail::errflags() = (detail::errflags()|(*flagp&excepts)) & (*flagp|~excepts); return 0; }

	/// Throw C++ exceptions based on set exception flags.
	/// This function manually throws a corresponding C++ exception if one of the specified flags is set, 
	/// no matter if automatic throwing (via [HALF_ERRHANDLING_THROW_...](\ref HALF_ERRHANDLING_THROW_INVALID)) is enabled or not.
	/// This function works even if [automatic exception flag handling](\ref HALF_ERRHANDLING_FLAGS) is disabled, 
	/// but in that case manual flag management is the only way to raise flags.
	/// \param excepts OR of exceptions to test
	/// \param msg error message to use for exception description
	/// \throw std::domain_error if `FE_INVALID` or `FE_DIVBYZERO` is selected and set
	/// \throw std::overflow_error if `FE_OVERFLOW` is selected and set
	/// \throw std::underflow_error if `FE_UNDERFLOW` is selected and set
	/// \throw std::range_error if `FE_INEXACT` is selected and set
	inline void fethrowexcept(int excepts, const char *msg = "") {
		excepts &= detail::errflags();
#if HALF_ERRHANDLING_THROWS
	#ifdef HALF_ERRHANDLING_THROW_INVALID
		if(excepts & FE_INVALID)
			throw std::domain_error(msg);
	#endif
	#ifdef HALF_ERRHANDLING_THROW_DIVBYZERO
		if(excepts & FE_DIVBYZERO)
			throw std::domain_error(msg);
	#endif
	#ifdef HALF_ERRHANDLING_THROW_OVERFLOW
		if(excepts & FE_OVERFLOW)
			throw std::overflow_error(msg);
	#endif
	#ifdef HALF_ERRHANDLING_THROW_UNDERFLOW
		if(excepts & FE_UNDERFLOW)
			throw std::underflow_error(msg);
	#endif
	#ifdef HALF_ERRHANDLING_THROW_INEXACT
		if(excepts & FE_INEXACT)
			throw std::range_error(msg);
	#endif
#else
		std::fprintf(stderr, "%s\n", msg);
		std::terminate();
#endif
	}
	/// \}
}


#undef HALF_UNUSED_NOERR
#undef constexpr_NOERR
#undef HALF_TWOS_COMPLEMENT_INT
#ifdef HALF_POP_WARNINGS
	#pragma warning(pop)
	#undef HALF_POP_WARNINGS
#endif