diff --git a/include/boost/decimal/decimal128_t.hpp b/include/boost/decimal/decimal128_t.hpp index 889664785..8b903b610 100644 --- a/include/boost/decimal/decimal128_t.hpp +++ b/include/boost/decimal/decimal128_t.hpp @@ -10,7 +10,7 @@ #include #include #include -#include "detail/int128.hpp" +#include #include #include #include @@ -122,18 +122,18 @@ BOOST_DECIMAL_EXPORT class decimal128_t final #endif - friend constexpr auto from_bits(int128::uint128_t rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto from_bits(int128::uint128_t rhs) noexcept -> decimal128_t; - constexpr auto unbiased_exponent() const noexcept -> std::uint64_t; - constexpr auto biased_exponent() const noexcept -> std::int32_t; - constexpr auto full_significand() const noexcept -> int128::uint128_t; - constexpr auto isneg() const noexcept -> bool; - constexpr auto to_components() const noexcept -> detail::decimal128_t_components; + BOOST_DECIMAL_CUDA_CONSTEXPR auto unbiased_exponent() const noexcept -> std::uint64_t; + BOOST_DECIMAL_CUDA_CONSTEXPR auto biased_exponent() const noexcept -> std::int32_t; + BOOST_DECIMAL_CUDA_CONSTEXPR auto full_significand() const noexcept -> int128::uint128_t; + BOOST_DECIMAL_CUDA_CONSTEXPR auto isneg() const noexcept -> bool; + BOOST_DECIMAL_CUDA_CONSTEXPR auto to_components() const noexcept -> detail::decimal128_t_components; // Allows direct editing of the exp template , bool> = true> - constexpr auto edit_exponent(T exp) noexcept -> void; - constexpr auto edit_sign(bool sign) noexcept -> void; + BOOST_DECIMAL_CUDA_CONSTEXPR auto edit_exponent(T exp) noexcept -> void; + BOOST_DECIMAL_CUDA_CONSTEXPR auto edit_sign(bool sign) noexcept -> void; // Attempts conversion to integral type: // If this is nan sets errno to EINVAL and returns 0 @@ -172,7 +172,7 @@ BOOST_DECIMAL_EXPORT class decimal128_t final -> std::enable_if_t<(detail::is_decimal_floating_point_v && detail::is_decimal_floating_point_v), bool>; - friend constexpr auto d128_div_impl(const decimal128_t& lhs, const decimal128_t& rhs, decimal128_t& q, decimal128_t& r) noexcept -> void; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto d128_div_impl(const decimal128_t& lhs, const decimal128_t& rhs, decimal128_t& q, decimal128_t& r) noexcept -> void; template friend constexpr auto ilogb(T d) noexcept @@ -182,7 +182,7 @@ BOOST_DECIMAL_EXPORT class decimal128_t final friend constexpr auto logb(T num) noexcept BOOST_DECIMAL_REQUIRES(detail::is_decimal_floating_point_v, T); - friend constexpr auto not_finite(decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto not_finite(decimal128_t rhs) noexcept -> bool; friend constexpr auto to_bid_d128(decimal128_t val) noexcept -> int128::uint128_t; @@ -223,7 +223,7 @@ BOOST_DECIMAL_EXPORT class decimal128_t final friend constexpr auto detail::write_payload(typename TargetDecimalType::significand_type payload_value) BOOST_DECIMAL_REQUIRES(detail::is_ieee_type_v, TargetDecimalType); - friend constexpr auto nan_conversion(const decimal128_t value) noexcept -> decimal128_t + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto nan_conversion(const decimal128_t value) noexcept -> decimal128_t { constexpr auto convert_nan_mask {detail::d128_snan_mask ^ detail::d128_nan_mask}; @@ -279,10 +279,10 @@ BOOST_DECIMAL_EXPORT class decimal128_t final #if !defined(BOOST_DECIMAL_ALLOW_IMPLICIT_CONVERSIONS) && !defined(BOOST_DECIMAL_ALLOW_IMPLICIT_INTEGER_CONVERSIONS) explicit #endif - constexpr decimal128_t(Integer val) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t(Integer val) noexcept; template - constexpr auto operator=(const Integer& val) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator=(const Integer& val) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&); // 3.2.5 initialization from coefficient and exponent: @@ -291,7 +291,7 @@ BOOST_DECIMAL_EXPORT class decimal128_t final #else template && detail::is_integral_v, bool> = true> #endif - constexpr decimal128_t(T1 coeff, T2 exp, detail::construction_sign_wrapper resultant_sign = construction_sign::positive) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t(T1 coeff, T2 exp, detail::construction_sign_wrapper resultant_sign = construction_sign::positive) noexcept; #ifdef BOOST_DECIMAL_HAS_CONCEPTS template @@ -305,9 +305,9 @@ BOOST_DECIMAL_EXPORT class decimal128_t final #else template && detail::is_integral_v, bool> = true> #endif - constexpr decimal128_t(T1 coeff, T2 exp) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t(T1 coeff, T2 exp) noexcept; - explicit constexpr decimal128_t(bool value) noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t(bool value) noexcept; #if !defined(BOOST_DECIMAL_DISABLE_CLIB) @@ -322,16 +322,16 @@ BOOST_DECIMAL_EXPORT class decimal128_t final #endif // 3.2.4.4 Conversion to integral type - explicit constexpr operator bool() const noexcept; - explicit constexpr operator int() const noexcept; - explicit constexpr operator unsigned() const noexcept; - explicit constexpr operator long() const noexcept; - explicit constexpr operator unsigned long() const noexcept; - explicit constexpr operator long long() const noexcept; - explicit constexpr operator unsigned long long() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator bool() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator int() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator unsigned() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator long() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator unsigned long() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator long long() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator unsigned long long() const noexcept; - explicit constexpr operator boost::int128::int128_t() const noexcept; - explicit constexpr operator boost::int128::uint128_t() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator boost::int128::int128_t() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator boost::int128::uint128_t() const noexcept; #ifdef BOOST_DECIMAL_HAS_INT128 explicit constexpr operator detail::builtin_int128_t() const noexcept; @@ -360,212 +360,212 @@ BOOST_DECIMAL_EXPORT class decimal128_t final #endif template && (detail::decimal_val_v > detail::decimal_val_v), bool> = true> - constexpr operator Decimal() const noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR operator Decimal() const noexcept; template && (detail::decimal_val_v <= detail::decimal_val_v), bool> = true> - explicit constexpr operator Decimal() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator Decimal() const noexcept; // cmath functions that are easier as friends - friend constexpr auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; - friend constexpr auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; - friend constexpr auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; - friend constexpr auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; - friend constexpr auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; - friend constexpr auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal128_t rhs) noexcept -> bool; // 3.2.7 unary arithmetic operators: - friend constexpr auto operator+(decimal128_t rhs) noexcept -> decimal128_t; - friend constexpr auto operator-(decimal128_t rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(decimal128_t rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal128_t rhs) noexcept -> decimal128_t; // 3.2.8 Binary arithmetic operators - friend constexpr auto operator+(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; template - friend constexpr auto operator+(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); template - friend constexpr auto operator+(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); - friend constexpr auto operator-(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; template - friend constexpr auto operator-(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); template - friend constexpr auto operator-(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); - friend constexpr auto operator*(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; template - friend constexpr auto operator*(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); template - friend constexpr auto operator*(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); - friend constexpr auto operator/(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; template - friend constexpr auto operator/(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); template - friend constexpr auto operator/(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t); - friend constexpr auto operator%(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator%(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; // 3.2.4.5 Increment and Decrement - constexpr auto operator++() noexcept -> decimal128_t&; - constexpr auto operator++(int) noexcept -> decimal128_t; // NOLINT : C++14 so constexpr implies const - constexpr auto operator--() noexcept -> decimal128_t&; - constexpr auto operator--(int) noexcept -> decimal128_t; // NOLINT : C++14 so constexpr implies const + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator++() noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator++(int) noexcept -> decimal128_t; // NOLINT : C++14 so constexpr implies const + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator--() noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator--(int) noexcept -> decimal128_t; // NOLINT : C++14 so constexpr implies const // 3.2.4.6 Compound Assignment - constexpr auto operator+=(decimal128_t rhs) noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+=(decimal128_t rhs) noexcept -> decimal128_t&; template - constexpr auto operator+=(Integer rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+=(Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&); template - constexpr auto operator+=(Decimal rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+=(Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&); - constexpr auto operator-=(decimal128_t rhs) noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-=(decimal128_t rhs) noexcept -> decimal128_t&; template - constexpr auto operator-=(Integer rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-=(Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&); template - constexpr auto operator-=(Decimal rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-=(Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&); - constexpr auto operator*=(decimal128_t rhs) noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*=(decimal128_t rhs) noexcept -> decimal128_t&; template - constexpr auto operator*=(Integer rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*=(Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&); template - constexpr auto operator*=(Decimal rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*=(Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&); - constexpr auto operator/=(decimal128_t rhs) noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/=(decimal128_t rhs) noexcept -> decimal128_t&; template - constexpr auto operator/=(Integer rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/=(Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&); template - constexpr auto operator/=(Decimal rhs) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/=(Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&); - constexpr auto operator%=(decimal128_t rhs) noexcept -> decimal128_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator%=(decimal128_t rhs) noexcept -> decimal128_t&; // 3.2.9 Comparison operators: // Equality - friend constexpr auto operator==(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; template - friend constexpr auto operator==(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator==(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Inequality - friend constexpr auto operator!=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; template - friend constexpr auto operator!=(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator!=(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Less - friend constexpr auto operator<(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; template - friend constexpr auto operator<(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator<(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Less equal - friend constexpr auto operator<=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; template - friend constexpr auto operator<=(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator<=(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Greater - friend constexpr auto operator>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; template - friend constexpr auto operator>(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator>(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Greater equal - friend constexpr auto operator>=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; template - friend constexpr auto operator>=(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator>=(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // C++20 spaceship #ifdef BOOST_DECIMAL_HAS_SPACESHIP_OPERATOR - friend constexpr auto operator<=>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> std::partial_ordering; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> std::partial_ordering; template - friend constexpr auto operator<=>(decimal128_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(decimal128_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering); template - friend constexpr auto operator<=>(Integer lhs, decimal128_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(Integer lhs, decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering); #endif // 3.6.4 Same Quantum - friend constexpr auto samequantumd128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto samequantumd128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool; // 3.6.5 Quantum exponent - friend constexpr auto quantexpd128(decimal128_t x) noexcept -> int; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto quantexpd128(decimal128_t x) noexcept -> int; // 3.6.6 Quantize - friend constexpr auto quantized128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto quantized128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t; // functions that need to be friends template friend constexpr auto frexp10(T num, int* expptr) noexcept -> typename T::significand_type; - friend constexpr auto copysignd128(decimal128_t mag, decimal128_t sgn) noexcept -> decimal128_t; - friend constexpr auto scalblnd128(decimal128_t num, long exp) noexcept -> decimal128_t; - friend constexpr auto scalbnd128(decimal128_t num, int exp) noexcept -> decimal128_t; - friend constexpr auto fmad128(decimal128_t x, decimal128_t y, decimal128_t z) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto copysignd128(decimal128_t mag, decimal128_t sgn) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto scalblnd128(decimal128_t num, long exp) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto scalbnd128(decimal128_t num, int exp) noexcept -> decimal128_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto fmad128(decimal128_t x, decimal128_t y, decimal128_t z) noexcept -> decimal128_t; }; #ifdef BOOST_DECIMAL_HAS_INT128 @@ -585,7 +585,7 @@ constexpr auto to_bits(const decimal128_t rhs) noexcept -> detail::builtin_uint1 #endif -constexpr auto from_bits(const int128::uint128_t rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto from_bits(const int128::uint128_t rhs) noexcept -> decimal128_t { decimal128_t result; result.bits_ = rhs; @@ -593,7 +593,7 @@ constexpr auto from_bits(const int128::uint128_t rhs) noexcept -> decimal128_t return result; } -constexpr auto decimal128_t::unbiased_exponent() const noexcept -> exponent_type +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::unbiased_exponent() const noexcept -> exponent_type { exponent_type expval {}; @@ -609,12 +609,12 @@ constexpr auto decimal128_t::unbiased_exponent() const noexcept -> exponent_type return expval; } -constexpr auto decimal128_t::biased_exponent() const noexcept -> biased_exponent_type +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::biased_exponent() const noexcept -> biased_exponent_type { return static_cast(unbiased_exponent()) - detail::bias_v; } -constexpr auto decimal128_t::full_significand() const noexcept -> significand_type +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::full_significand() const noexcept -> significand_type { significand_type significand {}; @@ -631,12 +631,12 @@ constexpr auto decimal128_t::full_significand() const noexcept -> significand_ty return significand; } -constexpr auto decimal128_t::isneg() const noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::isneg() const noexcept -> bool { return static_cast(bits_.high & detail::d128_sign_mask); } -constexpr auto decimal128_t::to_components() const noexcept -> detail::decimal128_t_components +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::to_components() const noexcept -> detail::decimal128_t_components { significand_type significand {}; exponent_type expval {}; @@ -660,12 +660,12 @@ constexpr auto decimal128_t::to_components() const noexcept -> detail::decimal12 template , bool>> -constexpr auto decimal128_t::edit_exponent(const T expval) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::edit_exponent(const T expval) noexcept -> void { *this = decimal128_t(this->full_significand(), expval, this->isneg()); } -constexpr auto decimal128_t::edit_sign(const bool sign) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::edit_sign(const bool sign) noexcept -> void { if (sign) { @@ -696,7 +696,7 @@ template #else template && detail::is_integral_v, bool>> #endif -constexpr decimal128_t::decimal128_t(T1 coeff, T2 exp, const detail::construction_sign_wrapper resultant_sign) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::decimal128_t(T1 coeff, T2 exp, const detail::construction_sign_wrapper resultant_sign) noexcept { const auto is_negative {static_cast(resultant_sign)}; bits_.high = is_negative ? detail::d128_sign_mask : UINT64_C(0); @@ -709,7 +709,10 @@ constexpr decimal128_t::decimal128_t(T1 coeff, T2 exp, const detail::constructio auto biased_exp {static_cast(exp + detail::bias_v)}; BOOST_DECIMAL_IF_CONSTEXPR (sizeof(T1) >= sizeof(significand_type)) { - if (coeff > detail::d128_max_significand_value || biased_exp < -(detail::precision_v - 1)) + // NVCC can't access uint128_t variables since it's a non-trivial struct. + // We need to call the related function and store the result instead. + constexpr auto max_sig {detail::impl::max_significand_v()}; + if (coeff > max_sig || biased_exp < -(detail::precision_v - 1)) { coeff_digits = detail::coefficient_rounding(coeff, exp, biased_exp, is_negative, detail::num_digits(coeff)); } @@ -854,9 +857,9 @@ template #else template && detail::is_integral_v, bool>> #endif -constexpr decimal128_t::decimal128_t(const T1 coeff, const T2 exp) noexcept : decimal128_t(detail::make_positive_unsigned(coeff), exp, coeff < 0) {} +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::decimal128_t(const T1 coeff, const T2 exp) noexcept : decimal128_t(detail::make_positive_unsigned(coeff), exp, coeff < 0) {} -constexpr decimal128_t::decimal128_t(const bool value) noexcept : decimal128_t(static_cast(value), 0, false) {} +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::decimal128_t(const bool value) noexcept : decimal128_t(static_cast(value), 0, false) {} #ifdef _MSC_VER @@ -904,15 +907,15 @@ class numeric_limits_impl128 static constexpr bool tinyness_before = true; // Member functions - static constexpr auto (min) () -> boost::decimal::decimal128_t { return {UINT32_C(1), min_exponent}; } - static constexpr auto (max) () -> boost::decimal::decimal128_t { return {d128_max_significand_value, max_exponent - digits + 1}; } - static constexpr auto lowest () -> boost::decimal::decimal128_t { return {d128_max_significand_value, max_exponent - digits + 1, construction_sign::negative}; } - static constexpr auto epsilon () -> boost::decimal::decimal128_t { return {UINT32_C(1), -digits + 1}; } - static constexpr auto round_error () -> boost::decimal::decimal128_t { return epsilon(); } - static constexpr auto infinity () -> boost::decimal::decimal128_t { return boost::decimal::from_bits(boost::decimal::detail::d128_inf_mask); } - static constexpr auto quiet_NaN () -> boost::decimal::decimal128_t { return boost::decimal::from_bits(boost::decimal::detail::d128_nan_mask); } - static constexpr auto signaling_NaN() -> boost::decimal::decimal128_t { return boost::decimal::from_bits(boost::decimal::detail::d128_snan_mask); } - static constexpr auto denorm_min () -> boost::decimal::decimal128_t { return {1, boost::decimal::detail::etiny_v}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto (min) () -> boost::decimal::decimal128_t { return {UINT32_C(1), min_exponent}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto (max) () -> boost::decimal::decimal128_t { return {d128_max_significand_value, max_exponent - digits + 1}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto lowest () -> boost::decimal::decimal128_t { return {d128_max_significand_value, max_exponent - digits + 1, construction_sign::negative}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto epsilon () -> boost::decimal::decimal128_t { return {UINT32_C(1), -digits + 1}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto round_error () -> boost::decimal::decimal128_t { return epsilon(); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto infinity () -> boost::decimal::decimal128_t { return boost::decimal::from_bits(boost::decimal::detail::d128_inf_mask); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto quiet_NaN () -> boost::decimal::decimal128_t { return boost::decimal::from_bits(boost::decimal::detail::d128_nan_mask); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto signaling_NaN() -> boost::decimal::decimal128_t { return boost::decimal::from_bits(boost::decimal::detail::d128_snan_mask); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto denorm_min () -> boost::decimal::decimal128_t { return {1, boost::decimal::detail::etiny_v}; } }; #if !defined(__cpp_inline_variables) || __cpp_inline_variables < 201606L @@ -1038,10 +1041,10 @@ template #else template , bool>> #endif -constexpr decimal128_t::decimal128_t(const Integer val) noexcept : decimal128_t{val, 0} {} +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::decimal128_t(const Integer val) noexcept : decimal128_t{val, 0} {} template -constexpr auto decimal128_t::operator=(const Integer& val) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator=(const Integer& val) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&) { using ConversionType = std::conditional_t::value, std::int32_t, Integer>; @@ -1059,48 +1062,48 @@ constexpr decimal128_t::decimal128_t(const Decimal val) noexcept *this = to_decimal(val); } -constexpr decimal128_t::operator bool() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator bool() const noexcept { constexpr decimal128_t zero {0, 0}; return *this != zero; } -constexpr decimal128_t::operator int() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator int() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator unsigned() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator unsigned() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator long() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator unsigned long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator unsigned long() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator long long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator long long() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator unsigned long long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator unsigned long long() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator boost::int128::int128_t() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator boost::int128::int128_t() const noexcept { return to_integral_128(*this); } -constexpr decimal128_t::operator boost::int128::uint128_t() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator boost::int128::uint128_t() const noexcept { return to_integral_128(*this); } @@ -1162,23 +1165,23 @@ constexpr decimal128_t::operator std::bfloat16_t() const noexcept #endif template && (detail::decimal_val_v > detail::decimal_val_v), bool>> -constexpr decimal128_t::operator Decimal() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator Decimal() const noexcept { return to_decimal(*this); } template && (detail::decimal_val_v <= detail::decimal_val_v), bool>> -constexpr decimal128_t::operator Decimal() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal128_t::operator Decimal() const noexcept { return to_decimal(*this); } -constexpr auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool { return rhs.bits_.high & detail::d128_sign_mask; } -constexpr auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_.high & detail::d128_nan_mask.high) == detail::d128_nan_mask.high; @@ -1188,7 +1191,7 @@ constexpr auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_ #endif } -constexpr auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_.high & detail::d128_nan_mask.high) == detail::d128_inf_mask.high; @@ -1198,7 +1201,7 @@ constexpr auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_ #endif } -constexpr auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_.high & detail::d128_snan_mask.high) == detail::d128_snan_mask.high; @@ -1208,7 +1211,7 @@ constexpr auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decim #endif } -constexpr auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH // Check for de-normals @@ -1226,7 +1229,7 @@ constexpr auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal1 #endif } -constexpr auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal128_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_.high & detail::d128_inf_mask.high) != detail::d128_inf_mask.high; @@ -1236,7 +1239,7 @@ constexpr auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal1 #endif } -BOOST_DECIMAL_FORCE_INLINE constexpr auto not_finite(const decimal128_t rhs) noexcept -> bool +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto not_finite(const decimal128_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_.high & detail::d128_inf_mask.high) == detail::d128_inf_mask.high; @@ -1246,57 +1249,57 @@ BOOST_DECIMAL_FORCE_INLINE constexpr auto not_finite(const decimal128_t rhs) noe #endif } -constexpr auto operator+(const decimal128_t rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal128_t rhs) noexcept -> decimal128_t { return rhs; } -constexpr auto operator-(decimal128_t rhs) noexcept-> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal128_t rhs) noexcept-> decimal128_t { rhs.bits_.high ^= detail::d128_sign_mask; return rhs; } -constexpr auto operator==(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { return equality_impl(lhs, rhs); } template -constexpr auto operator==(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return mixed_equality_impl(lhs, rhs); } template -constexpr auto operator==(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return mixed_equality_impl(rhs, lhs); } -constexpr auto operator!=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { return !(lhs == rhs); } template -constexpr auto operator!=(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return !(lhs == rhs); } template -constexpr auto operator!=(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return !(lhs == rhs); } -constexpr auto operator<(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1322,14 +1325,14 @@ constexpr auto operator<(const decimal128_t& lhs, const decimal128_t& rhs) noexc } template -constexpr auto operator<(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return less_impl(lhs, rhs); } template -constexpr auto operator<(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1342,7 +1345,7 @@ constexpr auto operator<(const Integer lhs, const decimal128_t rhs) noexcept return !less_impl(rhs, lhs) && lhs != rhs; } -constexpr auto operator<=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH if (isnan(lhs) || isnan(rhs)) @@ -1355,7 +1358,7 @@ constexpr auto operator<=(const decimal128_t& lhs, const decimal128_t& rhs) noex } template -constexpr auto operator<=(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1369,7 +1372,7 @@ constexpr auto operator<=(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator<=(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1382,26 +1385,26 @@ constexpr auto operator<=(const Integer lhs, const decimal128_t rhs) noexcept return !(rhs < lhs); } -constexpr auto operator>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { return rhs < lhs; } template -constexpr auto operator>(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return rhs < lhs; } template -constexpr auto operator>(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return rhs < lhs; } -constexpr auto operator>=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH if (isnan(lhs) || isnan(rhs)) @@ -1414,7 +1417,7 @@ constexpr auto operator>=(const decimal128_t& lhs, const decimal128_t& rhs) noex } template -constexpr auto operator>=(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1428,7 +1431,7 @@ constexpr auto operator>=(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator>=(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1443,7 +1446,7 @@ constexpr auto operator>=(const Integer lhs, const decimal128_t rhs) noexcept #ifdef BOOST_DECIMAL_HAS_SPACESHIP_OPERATOR -constexpr auto operator<=>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> std::partial_ordering +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> std::partial_ordering { if (lhs < rhs) { @@ -1462,7 +1465,7 @@ constexpr auto operator<=>(const decimal128_t& lhs, const decimal128_t& rhs) noe } template -constexpr auto operator<=>(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering) { if (lhs < rhs) @@ -1482,7 +1485,7 @@ constexpr auto operator<=>(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator<=>(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering) { if (lhs < rhs) @@ -1535,7 +1538,7 @@ std::ostream& operator<<( std::ostream& os, boost::decimal::detail::builtin_uint # pragma warning(disable: 4127) // If constexpr macro only works for C++17 and above #endif -constexpr auto d128_div_impl(const decimal128_t& lhs, const decimal128_t& rhs, decimal128_t& q, decimal128_t& r) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto d128_div_impl(const decimal128_t& lhs, const decimal128_t& rhs, decimal128_t& q, decimal128_t& r) noexcept -> void { #ifndef BOOST_DECIMAL_FAST_MATH // Check pre-conditions @@ -1653,7 +1656,7 @@ constexpr auto d128_div_impl(const decimal128_t& lhs, const decimal128_t& rhs, d # pragma warning(pop) #endif -constexpr auto operator+(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1676,7 +1679,7 @@ constexpr auto operator+(const decimal128_t& lhs, const decimal128_t& rhs) noexc } template -constexpr auto operator+(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { using exp_type = decimal128_t::biased_exponent_type; @@ -1701,14 +1704,14 @@ constexpr auto operator+(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator+(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { return rhs + lhs; } // NOLINTNEXTLINE: If subtraction is actually addition than use operator+ and vice versa -constexpr auto operator-(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1736,7 +1739,7 @@ constexpr auto operator-(const decimal128_t& lhs, const decimal128_t& rhs) noexc } template -constexpr auto operator-(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { using exp_type = decimal128_t::biased_exponent_type; @@ -1761,7 +1764,7 @@ constexpr auto operator-(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator-(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1779,7 +1782,7 @@ constexpr auto operator-(const Integer lhs, const decimal128_t rhs) noexcept return -rhs + lhs; } -constexpr auto operator*(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1818,7 +1821,7 @@ constexpr auto operator*(const decimal128_t& lhs, const decimal128_t& rhs) noexc } template -constexpr auto operator*(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { using exp_type = decimal128_t::biased_exponent_type; @@ -1857,13 +1860,13 @@ constexpr auto operator*(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator*(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { return rhs * lhs; } -constexpr auto operator/(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t { decimal128_t q {}; decimal128_t r {}; @@ -1873,7 +1876,7 @@ constexpr auto operator/(const decimal128_t& lhs, const decimal128_t& rhs) noexc } template -constexpr auto operator/(const decimal128_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const decimal128_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1915,7 +1918,7 @@ constexpr auto operator/(const decimal128_t lhs, const Integer rhs) noexcept } template -constexpr auto operator/(const Integer lhs, const decimal128_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const Integer lhs, const decimal128_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1951,7 +1954,7 @@ constexpr auto operator/(const Integer lhs, const decimal128_t rhs) noexcept return decimal128_t(q_components.sig, q_components.exp, q_components.sign); } -constexpr auto operator%(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator%(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t { decimal128_t q {}; decimal128_t r {}; @@ -1997,42 +2000,42 @@ constexpr auto operator%(const decimal128_t& lhs, const decimal128_t& rhs) noexc return r; } -constexpr auto decimal128_t::operator++() noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator++() noexcept -> decimal128_t& { constexpr decimal128_t one{1, 0}; *this = *this + one; return *this; } -constexpr auto decimal128_t::operator++(int) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator++(int) noexcept -> decimal128_t { const auto temp {*this}; ++(*this); return temp; } -constexpr auto decimal128_t::operator--() noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator--() noexcept -> decimal128_t& { constexpr decimal128_t one{1, 0}; *this = *this - one; return *this; } -constexpr auto decimal128_t::operator--(int) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator--(int) noexcept -> decimal128_t { const auto temp {*this}; --(*this); return temp; } -constexpr auto decimal128_t::operator+=(const decimal128_t rhs) noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator+=(const decimal128_t rhs) noexcept -> decimal128_t& { *this = *this + rhs; return *this; } template -constexpr auto decimal128_t::operator+=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator+=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&) { *this = *this + rhs; @@ -2040,21 +2043,21 @@ constexpr auto decimal128_t::operator+=(const Integer rhs) noexcept } template -constexpr auto decimal128_t::operator+=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator+=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&) { *this = *this + rhs; return *this; } -constexpr auto decimal128_t::operator-=(const decimal128_t rhs) noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator-=(const decimal128_t rhs) noexcept -> decimal128_t& { *this = *this - rhs; return *this; } template -constexpr auto decimal128_t::operator-=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator-=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&) { *this = *this - rhs; @@ -2062,21 +2065,21 @@ constexpr auto decimal128_t::operator-=(const Integer rhs) noexcept } template -constexpr auto decimal128_t::operator-=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator-=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&) { *this = *this - rhs; return *this; } -constexpr auto decimal128_t::operator*=(const decimal128_t rhs) noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator*=(const decimal128_t rhs) noexcept -> decimal128_t& { *this = *this * rhs; return *this; } template -constexpr auto decimal128_t::operator*=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator*=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&) { *this = *this * rhs; @@ -2084,21 +2087,21 @@ constexpr auto decimal128_t::operator*=(const Integer rhs) noexcept } template -constexpr auto decimal128_t::operator*=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator*=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&) { *this = *this * rhs; return *this; } -constexpr auto decimal128_t::operator/=(const decimal128_t rhs) noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator/=(const decimal128_t rhs) noexcept -> decimal128_t& { *this = *this / rhs; return *this; } template -constexpr auto decimal128_t::operator/=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator/=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal128_t&) { *this = *this / rhs; @@ -2106,14 +2109,14 @@ constexpr auto decimal128_t::operator/=(const Integer rhs) noexcept } template -constexpr auto decimal128_t::operator/=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator/=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal128_t&) { *this = *this / rhs; return *this; } -constexpr auto decimal128_t::operator%=(const decimal128_t rhs) noexcept -> decimal128_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal128_t::operator%=(const decimal128_t rhs) noexcept -> decimal128_t& { *this = *this % rhs; return *this; @@ -2124,7 +2127,7 @@ constexpr auto decimal128_t::operator%=(const decimal128_t rhs) noexcept -> deci // If both x and y are NaN, or infinity, they have the same quantum exponents; // if exactly one operand is infinity or exactly one operand is NaN, they do not have the same quantum exponents. // The samequantum functions raise no exception. -constexpr auto samequantumd128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto samequantumd128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH const auto lhs_fp {fpclassify(lhs)}; @@ -2146,7 +2149,7 @@ constexpr auto samequantumd128(const decimal128_t& lhs, const decimal128_t& rhs) // 3.6.5 // Effects: if x is finite, returns its quantum exponent. // Otherwise, a domain error occurs and INT_MIN is returned. -constexpr auto quantexpd128(const decimal128_t x) noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto quantexpd128(const decimal128_t x) noexcept -> int { #ifndef BOOST_DECIMAL_FAST_MATH if (!isfinite(x)) @@ -2169,7 +2172,7 @@ constexpr auto quantexpd128(const decimal128_t x) noexcept -> int // Otherwise, if only one operand is infinity, the "invalid" floating-point exception is raised and the result is NaN. // If both operands are infinity, the result is DEC_INFINITY, with the same sign as x, converted to the type of x. // The quantize functions do not signal underflow. -constexpr auto quantized128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto quantized128(const decimal128_t& lhs, const decimal128_t& rhs) noexcept -> decimal128_t { #ifndef BOOST_DECIMAL_FAST_MATH // Return the correct type of nan @@ -2196,13 +2199,13 @@ constexpr auto quantized128(const decimal128_t& lhs, const decimal128_t& rhs) no return {lhs.full_significand(), rhs.biased_exponent(), lhs.isneg()}; } -constexpr auto copysignd128(decimal128_t mag, const decimal128_t sgn) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto copysignd128(decimal128_t mag, const decimal128_t sgn) noexcept -> decimal128_t { mag.edit_sign(sgn.isneg()); return mag; } -constexpr auto scalblnd128(decimal128_t num, const long exp) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto scalblnd128(decimal128_t num, const long exp) noexcept -> decimal128_t { #ifndef BOOST_DECIMAL_FAST_MATH constexpr decimal128_t zero {0, 0}; @@ -2218,7 +2221,7 @@ constexpr auto scalblnd128(decimal128_t num, const long exp) noexcept -> decimal return num; } -constexpr auto scalbnd128(decimal128_t num, const int expval) noexcept -> decimal128_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto scalbnd128(decimal128_t num, const int expval) noexcept -> decimal128_t { return scalblnd128(num, static_cast(expval)); } diff --git a/include/boost/decimal/detail/add_impl.hpp b/include/boost/decimal/detail/add_impl.hpp index 4bd278367..d2603b04a 100644 --- a/include/boost/decimal/detail/add_impl.hpp +++ b/include/boost/decimal/detail/add_impl.hpp @@ -190,7 +190,7 @@ BOOST_DECIMAL_CUDA_CONSTEXPR auto add_impl(const T& lhs, const T& rhs) noexcept } template -constexpr auto d128_add_impl_new(const T& lhs, const T& rhs) noexcept -> ReturnType +BOOST_DECIMAL_CUDA_CONSTEXPR auto d128_add_impl_new(const T& lhs, const T& rhs) noexcept -> ReturnType { using promoted_sig_type = u256; diff --git a/include/boost/decimal/detail/div_impl.hpp b/include/boost/decimal/detail/div_impl.hpp index 877687d00..4bb9a5c8b 100644 --- a/include/boost/decimal/detail/div_impl.hpp +++ b/include/boost/decimal/detail/div_impl.hpp @@ -67,7 +67,7 @@ BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_generic_div_imp } template -constexpr auto d128_generic_div_impl(const T& lhs, const T& rhs, T& q) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto d128_generic_div_impl(const T& lhs, const T& rhs, T& q) noexcept -> void { bool sign {lhs.sign != rhs.sign}; diff --git a/include/boost/decimal/detail/i256.hpp b/include/boost/decimal/detail/i256.hpp index d15e4cb6f..790c45329 100644 --- a/include/boost/decimal/detail/i256.hpp +++ b/include/boost/decimal/detail/i256.hpp @@ -19,7 +19,7 @@ namespace detail { namespace impl { // This impl works regardless of intrinsics or otherwise -constexpr u256 u256_add_impl(const int128::uint128_t& lhs, const int128::uint128_t& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 u256_add_impl(const int128::uint128_t& lhs, const int128::uint128_t& rhs) noexcept { u256 result; std::uint64_t carry {}; @@ -39,7 +39,7 @@ constexpr u256 u256_add_impl(const int128::uint128_t& lhs, const int128::uint128 #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(BOOST_DECIMAL_ADD_CARRY) -constexpr u256 u256_add(const int128::uint128_t& lhs, const int128::uint128_t& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 u256_add(const int128::uint128_t& lhs, const int128::uint128_t& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -58,7 +58,7 @@ constexpr u256 u256_add(const int128::uint128_t& lhs, const int128::uint128_t& r #elif !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__GNUC__) && !defined(BOOST_DECIMAL_ADD_CARRY) -constexpr u256 u256_add(const int128::uint128_t& lhs, const int128::uint128_t& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 u256_add(const int128::uint128_t& lhs, const int128::uint128_t& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -79,7 +79,7 @@ constexpr u256 u256_add(const int128::uint128_t& lhs, const int128::uint128_t& r namespace impl { -constexpr std::uint64_t sub_borrow_u64(const std::uint64_t borrow_in, const std::uint64_t a, const std::uint64_t b, std::uint64_t& result) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR std::uint64_t sub_borrow_u64(const std::uint64_t borrow_in, const std::uint64_t a, const std::uint64_t b, std::uint64_t& result) noexcept { const auto diff {a - b}; const auto b1 {static_cast(a < b)}; @@ -89,7 +89,7 @@ constexpr std::uint64_t sub_borrow_u64(const std::uint64_t borrow_in, const std: return borrow_out; } -constexpr bool i256_sub_impl(const u256& a, const u256& b, u256& result) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool i256_sub_impl(const u256& a, const u256& b, u256& result) noexcept { if (a >= b) { @@ -114,7 +114,7 @@ constexpr bool i256_sub_impl(const u256& a, const u256& b, u256& result) noexcep #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(BOOST_DECIMAL_SUB_BORROW) -constexpr bool i256_sub(const u256& a, const u256& b, u256& res) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool i256_sub(const u256& a, const u256& b, u256& res) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -153,7 +153,7 @@ constexpr bool i256_sub(const u256& a, const u256& b, u256& res) noexcept #elif !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__GNUC__) && !defined(BOOST_DECIMAL_SUB_BORROW) && BOOST_DECIMAL_HAS_BUILTIN(__builtin_subcll) -constexpr bool i256_sub(const u256& a, const u256& b, u256& result) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool i256_sub(const u256& a, const u256& b, u256& result) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -205,7 +205,7 @@ inline std::uint64_t subcll(const std::uint64_t a, const std::uint64_t b, const } // namespace impl -constexpr bool i256_sub(const u256& a, const u256& b, u256& result) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool i256_sub(const u256& a, const u256& b, u256& result) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -240,7 +240,7 @@ constexpr bool i256_sub(const u256& a, const u256& b, u256& result) noexcept #else -constexpr bool i256_sub(const u256& a, const u256& b, u256& result) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool i256_sub(const u256& a, const u256& b, u256& result) noexcept { return impl::i256_sub_impl(a, b, result); } diff --git a/include/boost/decimal/detail/mul_impl.hpp b/include/boost/decimal/detail/mul_impl.hpp index 9763d4b57..bfa5b5e7e 100644 --- a/include/boost/decimal/detail/mul_impl.hpp +++ b/include/boost/decimal/detail/mul_impl.hpp @@ -148,7 +148,7 @@ BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_mul_impl(T lhs_ template BOOST_DECIMAL_FORCE_INLINE -constexpr auto d128_mul_impl(const T1& lhs_sig, const U1 lhs_exp, const bool lhs_sign, +BOOST_DECIMAL_CUDA_CONSTEXPR auto d128_mul_impl(const T1& lhs_sig, const U1 lhs_exp, const bool lhs_sign, const T2& rhs_sig, const U2 rhs_exp, const bool rhs_sign) noexcept -> ReturnType { using sig_type = T1; @@ -176,7 +176,7 @@ constexpr auto d128_mul_impl(const T1& lhs_sig, const U1 lhs_exp, const bool lhs template BOOST_DECIMAL_FORCE_INLINE -constexpr auto d128_fast_mul_impl(const T1& lhs_sig, const U1 lhs_exp, const bool lhs_sign, +BOOST_DECIMAL_CUDA_CONSTEXPR auto d128_fast_mul_impl(const T1& lhs_sig, const U1 lhs_exp, const bool lhs_sign, const T2& rhs_sig, const U2 rhs_exp, const bool rhs_sign) noexcept -> ReturnType { const bool sign {lhs_sign != rhs_sign}; @@ -194,7 +194,7 @@ constexpr auto d128_fast_mul_impl(const T1& lhs_sig, const U1 lhs_exp, const boo } template -BOOST_DECIMAL_FORCE_INLINE auto mul_impl(const decimal128_t_components& lhs, const decimal128_t_components& rhs) noexcept -> ReturnType +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto mul_impl(const decimal128_t_components& lhs, const decimal128_t_components& rhs) noexcept -> ReturnType { return d128_mul_impl(lhs.sig, lhs.exp, lhs.sign, rhs.sig, rhs.exp, rhs.sign); diff --git a/include/boost/decimal/detail/u256.hpp b/include/boost/decimal/detail/u256.hpp index 97aa75b9f..ce4c0224d 100644 --- a/include/boost/decimal/detail/u256.hpp +++ b/include/boost/decimal/detail/u256.hpp @@ -26,48 +26,48 @@ u256 std::uint64_t bytes[4] {}; // Constructors - constexpr u256() noexcept = default; - constexpr u256(const u256& other) noexcept = default; - constexpr u256(u256&& other) noexcept = default; - constexpr u256& operator=(const u256& other) noexcept = default; - constexpr u256& operator=(u256&& other) noexcept = default; + BOOST_DECIMAL_CUDA_CONSTEXPR u256() noexcept = default; + BOOST_DECIMAL_CUDA_CONSTEXPR u256(const u256& other) noexcept = default; + BOOST_DECIMAL_CUDA_CONSTEXPR u256(u256&& other) noexcept = default; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator=(const u256& other) noexcept = default; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator=(u256&& other) noexcept = default; - constexpr u256(std::uint64_t byte3, std::uint64_t byte2, std::uint64_t byte1, std::uint64_t byte0) noexcept; - constexpr u256(const int128::uint128_t x) noexcept { bytes[0] = x.low; bytes[1] = x.high; } - constexpr u256(const std::uint64_t x) noexcept { bytes[0] = x; } + BOOST_DECIMAL_CUDA_CONSTEXPR u256(std::uint64_t byte3, std::uint64_t byte2, std::uint64_t byte1, std::uint64_t byte0) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256(const int128::uint128_t x) noexcept { bytes[0] = x.low; bytes[1] = x.high; } + BOOST_DECIMAL_CUDA_CONSTEXPR u256(const std::uint64_t x) noexcept { bytes[0] = x; } - explicit constexpr operator std::uint64_t() const noexcept { return bytes[0]; } + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator std::uint64_t() const noexcept { return bytes[0]; } template - explicit constexpr operator std::enable_if_t< + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator std::enable_if_t< !std::is_same::value, T>() const noexcept { return static_cast(bytes[0]); } // Conversion to/from int128::uint128_t - constexpr u256(const int128::uint128_t& high_, const int128::uint128_t& low_) noexcept; - explicit constexpr operator int128::uint128_t() const noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256(const int128::uint128_t& high_, const int128::uint128_t& low_) noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator int128::uint128_t() const noexcept; BOOST_DECIMAL_CUDA_CONSTEXPR std::uint64_t operator[](std::size_t i) const noexcept; BOOST_DECIMAL_CUDA_CONSTEXPR std::uint64_t& operator[](std::size_t i) noexcept; // Compound operators - constexpr u256& operator<<=(int amount) noexcept; - constexpr u256& operator>>=(int amount) noexcept; - constexpr u256& operator|=(const u256& rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator<<=(int amount) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator>>=(int amount) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator|=(const u256& rhs) noexcept; - constexpr u256& operator*=(const u256& rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator*=(const u256& rhs) noexcept; - constexpr u256& operator/=(const u256& rhs) noexcept; - constexpr u256& operator/=(const int128::uint128_t& rhs) noexcept; - constexpr u256& operator/=(std::uint64_t rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator/=(const u256& rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator/=(const int128::uint128_t& rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator/=(std::uint64_t rhs) noexcept; - constexpr u256& operator%=(const u256& rhs) noexcept; - constexpr u256& operator%=(std::uint64_t rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator%=(const u256& rhs) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator%=(std::uint64_t rhs) noexcept; - constexpr u256& operator++() noexcept; - constexpr u256& operator++(int) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator++() noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR u256& operator++(int) noexcept; }; } // namespace detail @@ -78,7 +78,7 @@ namespace boost { namespace decimal { namespace detail { -constexpr u256::u256(const std::uint64_t byte3, const std::uint64_t byte2, const std::uint64_t byte1, const std::uint64_t byte0) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256::u256(const std::uint64_t byte3, const std::uint64_t byte2, const std::uint64_t byte1, const std::uint64_t byte0) noexcept { bytes[0] = byte0; bytes[1] = byte1; @@ -86,7 +86,7 @@ constexpr u256::u256(const std::uint64_t byte3, const std::uint64_t byte2, const bytes[3] = byte3; } -constexpr u256::u256(const int128::uint128_t& high_, const int128::uint128_t& low_) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256::u256(const int128::uint128_t& high_, const int128::uint128_t& low_) noexcept { bytes[0] = low_.low; bytes[1] = low_.high; @@ -94,7 +94,7 @@ constexpr u256::u256(const int128::uint128_t& high_, const int128::uint128_t& lo bytes[3] = high_.high; } -constexpr u256::operator int128::uint128_t() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256::operator int128::uint128_t() const noexcept { return int128::uint128_t {bytes[1], bytes[0]}; } @@ -117,12 +117,12 @@ BOOST_DECIMAL_CUDA_CONSTEXPR std::uint64_t& u256::operator[](const std::size_t i namespace impl { -BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_equality_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR bool basic_equality_impl(const u256& lhs, const u256& rhs) noexcept { return lhs[0] == rhs[0] && lhs[1] == rhs[1] && lhs[2] == rhs[2] && lhs[3] == rhs[3]; } -BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_inequality_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR bool basic_inequality_impl(const u256& lhs, const u256& rhs) noexcept { return lhs[0] != rhs[0] || lhs[1] != rhs[1] || lhs[2] != rhs[2] || lhs[3] != rhs[3]; } @@ -131,7 +131,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_inequality_impl(const u256& lhs, #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__AVX2__) -constexpr bool operator==(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator==(const u256& lhs, const u256& rhs) noexcept { // Start comp from low word since they will most likely be filled if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) @@ -153,7 +153,7 @@ constexpr bool operator==(const u256& lhs, const u256& rhs) noexcept #else -constexpr bool operator==(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator==(const u256& lhs, const u256& rhs) noexcept { return impl::basic_equality_impl(lhs, rhs); } @@ -166,7 +166,7 @@ constexpr bool operator==(const u256& lhs, const u256& rhs) noexcept #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__AVX2__) -constexpr bool operator!=(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator!=(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -187,7 +187,7 @@ constexpr bool operator!=(const u256& lhs, const u256& rhs) noexcept #else -constexpr bool operator!=(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator!=(const u256& lhs, const u256& rhs) noexcept { return impl::basic_inequality_impl(lhs, rhs); } @@ -200,7 +200,7 @@ constexpr bool operator!=(const u256& lhs, const u256& rhs) noexcept namespace impl { -BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_lt_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR bool basic_lt_impl(const u256& lhs, const u256& rhs) noexcept { if (lhs[3] != rhs[3]) { @@ -222,7 +222,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_lt_impl(const u256& lhs, const u #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__AVX2__) -constexpr bool operator<(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -259,29 +259,29 @@ constexpr bool operator<(const u256& lhs, const u256& rhs) noexcept #else -constexpr bool operator<(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<(const u256& lhs, const u256& rhs) noexcept { return impl::basic_lt_impl(lhs, rhs); } #endif -constexpr bool operator<(const u256& lhs, const int128::uint128_t& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<(const u256& lhs, const int128::uint128_t& rhs) noexcept { return lhs[3] == 0U && lhs[2] == 0U && int128::uint128_t{lhs[1], lhs[0]} < rhs; } -constexpr bool operator<(const int128::uint128_t& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<(const int128::uint128_t& lhs, const u256& rhs) noexcept { return rhs[3] == 0U && rhs[2] == 0U && lhs < int128::uint128_t{rhs[1], rhs[0]}; } -constexpr bool operator<(const u256& lhs, const std::uint64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<(const u256& lhs, const std::uint64_t rhs) noexcept { return lhs[3] == 0 && lhs[2] == 0 && lhs[1] == 0 && lhs[0] < rhs; } -constexpr bool operator<(const std::uint64_t lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<(const std::uint64_t lhs, const u256& rhs) noexcept { return rhs[3] == 0 && rhs[2] == 0 && rhs[1] == 0 && lhs < rhs[0]; } @@ -292,7 +292,7 @@ constexpr bool operator<(const std::uint64_t lhs, const u256& rhs) noexcept namespace impl { -BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_le_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR bool basic_le_impl(const u256& lhs, const u256& rhs) noexcept { return !(rhs < lhs); } @@ -301,7 +301,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr bool basic_le_impl(const u256& lhs, const u #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__AVX2__) -constexpr bool operator<=(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<=(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -342,19 +342,19 @@ constexpr bool operator<=(const u256& lhs, const u256& rhs) noexcept #else -constexpr bool operator<=(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<=(const u256& lhs, const u256& rhs) noexcept { return impl::basic_le_impl(lhs, rhs); } #endif -constexpr bool operator<=(const u256& lhs, const std::uint64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<=(const u256& lhs, const std::uint64_t rhs) noexcept { return lhs[3] == 0 && lhs[2] == 0 && lhs[1] == 0 && lhs[0] <= rhs; } -constexpr bool operator<=(const std::uint64_t lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator<=(const std::uint64_t lhs, const u256& rhs) noexcept { return rhs[3] == 0 && rhs[2] == 0 && rhs[1] == 0 && lhs <= rhs[0]; } @@ -363,17 +363,17 @@ constexpr bool operator<=(const std::uint64_t lhs, const u256& rhs) noexcept // Greater Than Operator //===================================== -constexpr bool operator>(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator>(const u256& lhs, const u256& rhs) noexcept { return rhs < lhs; } -constexpr bool operator>(const u256& lhs, const int128::uint128_t& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator>(const u256& lhs, const int128::uint128_t& rhs) noexcept { return lhs[3] > 0U || lhs[2] > 0U || int128::uint128_t{lhs[1], lhs[0]} > rhs; } -constexpr bool operator>(const u256& lhs, const std::uint64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator>(const u256& lhs, const std::uint64_t rhs) noexcept { return lhs[3] != 0U || lhs[2] != 0U || lhs[1] != 0U || lhs[0] > rhs; } @@ -382,7 +382,7 @@ constexpr bool operator>(const u256& lhs, const std::uint64_t rhs) noexcept // Greater Equal Operator //===================================== -constexpr bool operator>=(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR bool operator>=(const u256& lhs, const u256& rhs) noexcept { return !(lhs < rhs); } @@ -391,7 +391,7 @@ constexpr bool operator>=(const u256& lhs, const u256& rhs) noexcept // Left Shift Operators //===================================== -constexpr u256 operator<<(const u256& lhs, const int shift) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator<<(const u256& lhs, const int shift) noexcept { u256 result {}; @@ -428,7 +428,7 @@ constexpr u256 operator<<(const u256& lhs, const int shift) noexcept return result; } -constexpr u256& u256::operator<<=(const int amount) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator<<=(const int amount) noexcept { *this = *this << amount; return *this; @@ -438,7 +438,7 @@ constexpr u256& u256::operator<<=(const int amount) noexcept // Right Shift Operators //===================================== -constexpr u256 operator>>(const u256& lhs, const int shift) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator>>(const u256& lhs, const int shift) noexcept { u256 result {}; @@ -477,7 +477,7 @@ constexpr u256 operator>>(const u256& lhs, const int shift) noexcept return result; } -constexpr u256& u256::operator>>=(const int amount) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator>>=(const int amount) noexcept { *this = *this >> amount; return *this; @@ -489,7 +489,7 @@ constexpr u256& u256::operator>>=(const int amount) noexcept namespace impl { -BOOST_DECIMAL_FORCE_INLINE constexpr u256 basic_or_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 basic_or_impl(const u256& lhs, const u256& rhs) noexcept { u256 result; @@ -505,7 +505,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 basic_or_impl(const u256& lhs, const u #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__AVX2__) -constexpr u256 operator|(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator|(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -531,7 +531,7 @@ constexpr u256 operator|(const u256& lhs, const u256& rhs) noexcept #elif !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(BOOST_DECIMAL_HAS_ARM_INTRINSICS) -constexpr u256 operator|(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator|(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -561,14 +561,14 @@ constexpr u256 operator|(const u256& lhs, const u256& rhs) noexcept #else -constexpr u256 operator|(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator|(const u256& lhs, const u256& rhs) noexcept { return impl::basic_or_impl(lhs, rhs); } #endif -constexpr u256& u256::operator|=(const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator|=(const u256& rhs) noexcept { *this = *this | rhs; return *this; @@ -580,7 +580,7 @@ constexpr u256& u256::operator|=(const u256& rhs) noexcept namespace impl { -BOOST_DECIMAL_FORCE_INLINE constexpr u256 basic_and_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 basic_and_impl(const u256& lhs, const u256& rhs) noexcept { u256 result; @@ -596,7 +596,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 basic_and_impl(const u256& lhs, const #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(__AVX2__) -constexpr u256 operator&(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator&(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -622,7 +622,7 @@ constexpr u256 operator&(const u256& lhs, const u256& rhs) noexcept #elif !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(BOOST_DECIMAL_HAS_ARM_INTRINSICS) -constexpr u256 operator&(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator&(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -652,7 +652,7 @@ constexpr u256 operator&(const u256& lhs, const u256& rhs) noexcept #else -constexpr u256 operator&(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator&(const u256& lhs, const u256& rhs) noexcept { return impl::basic_and_impl(lhs, rhs); } @@ -665,7 +665,7 @@ constexpr u256 operator&(const u256& lhs, const u256& rhs) noexcept namespace impl { -constexpr u256 basic_add_impl(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 basic_add_impl(const u256& lhs, const u256& rhs) noexcept { u256 result; std::uint64_t carry {}; @@ -691,7 +691,7 @@ constexpr u256 basic_add_impl(const u256& lhs, const u256& rhs) noexcept #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && defined(BOOST_DECIMAL_ADD_CARRY) -constexpr u256 operator+(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator+(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -725,7 +725,7 @@ inline bool add_carry_u64(const bool carry_in, const std::uint64_t a, const std: } // namespace impl -constexpr u256 operator+(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator+(const u256& lhs, const u256& rhs) noexcept { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(lhs)) { @@ -746,20 +746,20 @@ constexpr u256 operator+(const u256& lhs, const u256& rhs) noexcept #else -constexpr u256 operator+(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator+(const u256& lhs, const u256& rhs) noexcept { return impl::basic_add_impl(lhs, rhs); } #endif -constexpr u256& u256::operator++() noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator++() noexcept { *this = *this + static_cast(1); return *this; } -constexpr u256& u256::operator++(int) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator++(int) noexcept { *this = *this + static_cast(1); return *this; @@ -778,7 +778,7 @@ namespace impl { #endif template -BOOST_DECIMAL_FORCE_INLINE constexpr u256 from_words(const std::uint32_t (&words)[word_size]) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 from_words(const std::uint32_t (&words)[word_size]) noexcept { static_assert(word_size >= 8, "Not enough words to convert to u256"); @@ -806,7 +806,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 from_words(const std::uint32_t (&words #endif template -constexpr u256 knuth_mulitply(const std::uint32_t (&u)[u_size], +BOOST_DECIMAL_CUDA_CONSTEXPR u256 knuth_mulitply(const std::uint32_t (&u)[u_size], const std::uint32_t (&v)[v_size]) noexcept { std::uint32_t w[u_size + v_size] {}; @@ -838,7 +838,7 @@ constexpr u256 knuth_mulitply(const std::uint32_t (&u)[u_size], return from_words(w); } -constexpr void to_words(const u256& x, std::uint32_t (&words)[8]) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR void to_words(const u256& x, std::uint32_t (&words)[8]) noexcept { #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && !BOOST_DECIMAL_ENDIAN_BIG_BYTE if (!BOOST_DECIMAL_DETAIL_INT128_IS_CONSTANT_EVALUATED(x)) @@ -861,7 +861,7 @@ constexpr void to_words(const u256& x, std::uint32_t (&words)[8]) noexcept } template -BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_mul(const u256& lhs, const UnsignedInteger& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 default_mul(const u256& lhs, const UnsignedInteger& rhs) noexcept { using boost::decimal::detail::impl::to_words; using boost::int128::detail::to_words; @@ -879,24 +879,24 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_mul(const u256& lhs, const Uns } // namespace impl -constexpr u256 operator*(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator*(const u256& lhs, const u256& rhs) noexcept { return impl::default_mul(lhs, rhs); } template -constexpr u256 operator*(const u256& lhs, const UnsignedInteger rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator*(const u256& lhs, const UnsignedInteger rhs) noexcept { return impl::default_mul(lhs, rhs); } template -constexpr u256 operator*(const UnsignedInteger lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator*(const UnsignedInteger lhs, const u256& rhs) noexcept { return impl::default_mul(rhs, lhs); } -constexpr u256 umul256(const int128::uint128_t& a, const int128::uint128_t& b) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 umul256(const int128::uint128_t& a, const int128::uint128_t& b) noexcept { u256 result{}; @@ -926,7 +926,7 @@ constexpr u256 umul256(const int128::uint128_t& a, const int128::uint128_t& b) n // 128×64→256 multiplication (SoftFloat-style lightweight primitive) // Used when rhs is 64-bit (e.g. r_scaled from approx_recip_sqrt64) // Explicit uint128_t cast ensures 64×64→128 widening (a.low*b otherwise returns uint64_t on some platforms) -constexpr u256 mul128By64(const int128::uint128_t& a, const std::uint64_t b) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 mul128By64(const int128::uint128_t& a, const std::uint64_t b) noexcept { const int128::uint128_t p0 = int128::uint128_t{a.low} * b; // 64×64→128 const int128::uint128_t p1 = int128::uint128_t{a.high} * b; // 64×64→128 @@ -943,7 +943,7 @@ constexpr u256 mul128By64(const int128::uint128_t& a, const std::uint64_t b) noe return result; } -constexpr u256& u256::operator*=(const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator*=(const u256& rhs) noexcept { *this = *this * rhs; return *this; @@ -955,7 +955,7 @@ constexpr u256& u256::operator*=(const u256& rhs) noexcept namespace impl { -constexpr std::size_t div_to_words(const u256& x, std::uint32_t (&words)[8]) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR std::size_t div_to_words(const u256& x, std::uint32_t (&words)[8]) noexcept { #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && !BOOST_DECIMAL_ENDIAN_BIG_BYTE if (!BOOST_DECIMAL_IS_CONSTANT_EVALUATED(x)) @@ -984,7 +984,7 @@ constexpr std::size_t div_to_words(const u256& x, std::uint32_t (&words)[8]) noe return word_count; } -constexpr std::size_t div_to_words(const boost::int128::uint128_t& x, std::uint32_t (&words)[8]) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR std::size_t div_to_words(const boost::int128::uint128_t& x, std::uint32_t (&words)[8]) noexcept { #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && !BOOST_DECIMAL_ENDIAN_BIG_BYTE if (!BOOST_DECIMAL_IS_CONSTANT_EVALUATED(x)) @@ -1011,7 +1011,7 @@ constexpr std::size_t div_to_words(const boost::int128::uint128_t& x, std::uint3 return word_count; } -constexpr std::size_t div_to_words(const std::uint64_t x, std::uint32_t (&words)[2]) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR std::size_t div_to_words(const std::uint64_t x, std::uint32_t (&words)[2]) noexcept { #if !defined(BOOST_DECIMAL_NO_CONSTEVAL_DETECTION) && !BOOST_DECIMAL_ENDIAN_BIG_BYTE if (!BOOST_DECIMAL_IS_CONSTANT_EVALUATED(x)) @@ -1036,7 +1036,7 @@ constexpr std::size_t div_to_words(const std::uint64_t x, std::uint32_t (&words) return word_count; } -BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_div(const u256& lhs, const std::uint64_t rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 default_div(const u256& lhs, const std::uint64_t rhs) noexcept { u256 quotient; @@ -1059,7 +1059,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_div(const u256& lhs, const std } template -BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_div(const u256& lhs, const UnsignedInteger& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 default_div(const u256& lhs, const UnsignedInteger& rhs) noexcept { if (rhs <= UINT64_MAX) { @@ -1091,7 +1091,7 @@ struct u256_divmod_result }; template -BOOST_DECIMAL_FORCE_INLINE constexpr auto div_mod(const u256& lhs, const UnsignedInteger& rhs) noexcept -> u256_divmod_result +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto div_mod(const u256& lhs, const UnsignedInteger& rhs) noexcept -> u256_divmod_result { std::uint32_t u[8] {}; std::uint32_t v[8] {}; @@ -1127,30 +1127,30 @@ BOOST_DECIMAL_FORCE_INLINE constexpr auto div_mod(const u256& lhs, const Unsigne } // namespace impl -constexpr u256 operator/(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator/(const u256& lhs, const u256& rhs) noexcept { return impl::default_div(lhs, rhs); } template -constexpr u256 operator/(const u256& lhs, const UnsignedInteger rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator/(const u256& lhs, const UnsignedInteger rhs) noexcept { return impl::default_div(lhs, rhs); } -constexpr u256& u256::operator/=(const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator/=(const u256& rhs) noexcept { *this = *this / rhs; return *this; } -constexpr u256& u256::operator/=(const int128::uint128_t& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator/=(const int128::uint128_t& rhs) noexcept { *this = *this / rhs; return *this; } -constexpr u256& u256::operator/=(const std::uint64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator/=(const std::uint64_t rhs) noexcept { *this = impl::default_div(*this, rhs); return *this; @@ -1162,7 +1162,7 @@ constexpr u256& u256::operator/=(const std::uint64_t rhs) noexcept namespace impl { -BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_mod(const u256& lhs, const std::uint64_t rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 default_mod(const u256& lhs, const std::uint64_t rhs) noexcept { u256 quotient; @@ -1186,7 +1186,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_mod(const u256& lhs, const std } template -BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_mod(const u256& lhs, const UnsignedInteger& rhs) noexcept +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR u256 default_mod(const u256& lhs, const UnsignedInteger& rhs) noexcept { if (rhs <= UINT64_MAX) { @@ -1207,24 +1207,24 @@ BOOST_DECIMAL_FORCE_INLINE constexpr u256 default_mod(const u256& lhs, const Uns } // namespace impl -constexpr u256 operator%(const u256& lhs, const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator%(const u256& lhs, const u256& rhs) noexcept { return impl::default_mod(lhs, rhs); } template -constexpr u256 operator%(const u256& lhs, const UnsignedInteger rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256 operator%(const u256& lhs, const UnsignedInteger rhs) noexcept { return impl::default_mod(lhs, rhs); } -constexpr u256& u256::operator%=(const u256& rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator%=(const u256& rhs) noexcept { *this = *this % rhs; return *this; } -constexpr u256& u256::operator%=(const std::uint64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR u256& u256::operator%=(const std::uint64_t rhs) noexcept { *this = impl::default_mod(*this, rhs); return *this; diff --git a/test/cuda_jamfile b/test/cuda_jamfile index c8122f67a..4db759979 100644 --- a/test/cuda_jamfile +++ b/test/cuda_jamfile @@ -15,3 +15,9 @@ run test_cuda_decimal64_t_add.cu ; run test_cuda_decimal64_t_sub.cu ; run test_cuda_decimal64_t_mul.cu ; run test_cuda_decimal64_t_div.cu ; + +run test_cuda_decimal128_t_construct.cu ; +run test_cuda_decimal128_t_add.cu ; +run test_cuda_decimal128_t_sub.cu ; +run test_cuda_decimal128_t_mul.cu ; +run test_cuda_decimal128_t_div.cu ; diff --git a/test/test_cuda_decimal128_t_add.cu b/test/test_cuda_decimal128_t_add.cu new file mode 100644 index 000000000..6233245dc --- /dev/null +++ b/test/test_cuda_decimal128_t_add.cu @@ -0,0 +1,77 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" +#include "stopwatch.hpp" + +#include + +using test_type = boost::decimal::decimal128_t; + +__global__ void cuda_test(const test_type *in, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = in[i] + in[i]; + } +} + +int main(void) +{ + std::mt19937_64 rng{42}; + + int numElements = 50000; + std::cout << "[Vector operation on " << numElements << " elements]" << std::endl; + + cuda_managed_ptr input_vector(numElements); + cuda_managed_ptr output_vector(numElements); + + std::uniform_int_distribution dist{1, 4999}; + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = test_type(dist(rng)); + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl; + + watch w; + + cudaDeviceSetLimit(cudaLimitStackSize, 8192); + cuda_test<<>>(input_vector.get(), output_vector.get(), numElements); + cudaDeviceSynchronize(); + + std::cout << "CUDA kernel done in: " << w.elapsed() << "s" << std::endl; + + std::vector results; + results.reserve(numElements); + w.reset(); + for (int i = 0; i < numElements; ++i) + { + results.push_back(input_vector[i] + input_vector[i]); + } + double t = w.elapsed(); + + for (int i = 0; i < numElements; ++i) + { + if (output_vector[i] != results[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl; + std::cout << "Done\n"; + + return 0; +} diff --git a/test/test_cuda_decimal128_t_construct.cu b/test/test_cuda_decimal128_t_construct.cu new file mode 100644 index 000000000..bb45fee4c --- /dev/null +++ b/test/test_cuda_decimal128_t_construct.cu @@ -0,0 +1,77 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" +#include "stopwatch.hpp" + +#include + +using test_type = boost::decimal::decimal128_t; + +__global__ void cuda_test(const int *in, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = test_type(in[i]); + } +} + +int main(void) +{ + std::mt19937_64 rng{42}; + + int numElements = 50000; + std::cout << "[Vector operation on " << numElements << " elements]" << std::endl; + + cuda_managed_ptr input_vector(numElements); + cuda_managed_ptr output_vector(numElements); + + std::uniform_int_distribution dist{-9999999, 9999999}; + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = dist(rng); + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl; + + watch w; + + cudaDeviceSetLimit(cudaLimitStackSize, 2048); + cuda_test<<>>(input_vector.get(), output_vector.get(), numElements); + cudaDeviceSynchronize(); + + std::cout << "CUDA kernel done in: " << w.elapsed() << "s" << std::endl; + + std::vector results; + results.reserve(numElements); + w.reset(); + for (int i = 0; i < numElements; ++i) + { + results.push_back(test_type(input_vector[i])); + } + double t = w.elapsed(); + + for (int i = 0; i < numElements; ++i) + { + if (output_vector[i] != results[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl; + std::cout << "Done\n"; + + return 0; +} diff --git a/test/test_cuda_decimal128_t_div.cu b/test/test_cuda_decimal128_t_div.cu new file mode 100644 index 000000000..a82c59492 --- /dev/null +++ b/test/test_cuda_decimal128_t_div.cu @@ -0,0 +1,80 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" +#include "stopwatch.hpp" + +#include + +using test_type = boost::decimal::decimal128_t; + +__global__ void cuda_test(const test_type *in, const test_type *in2, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = in[i] / in2[i]; + } +} + +int main(void) +{ + std::mt19937_64 rng{42}; + + int numElements = 50000; + std::cout << "[Vector operation on " << numElements << " elements]" << std::endl; + + cuda_managed_ptr input_vector(numElements); + cuda_managed_ptr input_vector2(numElements); + cuda_managed_ptr output_vector(numElements); + + std::uniform_int_distribution dist{1, 9999}; + std::uniform_int_distribution dist2{1, 999}; + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = test_type(dist(rng)); + input_vector2[i] = test_type(dist2(rng)); + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl; + + watch w; + + cudaDeviceSetLimit(cudaLimitStackSize, 8192); + cuda_test<<>>(input_vector.get(), input_vector2.get(), output_vector.get(), numElements); + cudaDeviceSynchronize(); + + std::cout << "CUDA kernel done in: " << w.elapsed() << "s" << std::endl; + + std::vector results; + results.reserve(numElements); + w.reset(); + for (int i = 0; i < numElements; ++i) + { + results.push_back(input_vector[i] / input_vector2[i]); + } + double t = w.elapsed(); + + for (int i = 0; i < numElements; ++i) + { + if (output_vector[i] != results[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl; + std::cout << "Done\n"; + + return 0; +} diff --git a/test/test_cuda_decimal128_t_mul.cu b/test/test_cuda_decimal128_t_mul.cu new file mode 100644 index 000000000..89fcc1e6f --- /dev/null +++ b/test/test_cuda_decimal128_t_mul.cu @@ -0,0 +1,79 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" +#include "stopwatch.hpp" + +#include + +using test_type = boost::decimal::decimal128_t; + +__global__ void cuda_test(const test_type *in, const test_type *in2, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = in[i] * in2[i]; + } +} + +int main(void) +{ + std::mt19937_64 rng{42}; + + int numElements = 50000; + std::cout << "[Vector operation on " << numElements << " elements]" << std::endl; + + cuda_managed_ptr input_vector(numElements); + cuda_managed_ptr input_vector2(numElements); + cuda_managed_ptr output_vector(numElements); + + std::uniform_int_distribution dist{1, 99}; + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = test_type(dist(rng)); + input_vector2[i] = test_type(dist(rng)); + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl; + + watch w; + + cudaDeviceSetLimit(cudaLimitStackSize, 8192); + cuda_test<<>>(input_vector.get(), input_vector2.get(), output_vector.get(), numElements); + cudaDeviceSynchronize(); + + std::cout << "CUDA kernel done in: " << w.elapsed() << "s" << std::endl; + + std::vector results; + results.reserve(numElements); + w.reset(); + for (int i = 0; i < numElements; ++i) + { + results.push_back(input_vector[i] * input_vector2[i]); + } + double t = w.elapsed(); + + for (int i = 0; i < numElements; ++i) + { + if (output_vector[i] != results[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl; + std::cout << "Done\n"; + + return 0; +} diff --git a/test/test_cuda_decimal128_t_sub.cu b/test/test_cuda_decimal128_t_sub.cu new file mode 100644 index 000000000..d9c5f9238 --- /dev/null +++ b/test/test_cuda_decimal128_t_sub.cu @@ -0,0 +1,79 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" +#include "stopwatch.hpp" + +#include + +using test_type = boost::decimal::decimal128_t; + +__global__ void cuda_test(const test_type *in, const test_type *in2, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = in[i] - in2[i]; + } +} + +int main(void) +{ + std::mt19937_64 rng{42}; + + int numElements = 50000; + std::cout << "[Vector operation on " << numElements << " elements]" << std::endl; + + cuda_managed_ptr input_vector(numElements); + cuda_managed_ptr input_vector2(numElements); + cuda_managed_ptr output_vector(numElements); + + std::uniform_int_distribution dist{1, 9999}; + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = test_type(dist(rng)); + input_vector2[i] = test_type(dist(rng)); + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl; + + watch w; + + cudaDeviceSetLimit(cudaLimitStackSize, 8192); + cuda_test<<>>(input_vector.get(), input_vector2.get(), output_vector.get(), numElements); + cudaDeviceSynchronize(); + + std::cout << "CUDA kernel done in: " << w.elapsed() << "s" << std::endl; + + std::vector results; + results.reserve(numElements); + w.reset(); + for (int i = 0; i < numElements; ++i) + { + results.push_back(input_vector[i] - input_vector2[i]); + } + double t = w.elapsed(); + + for (int i = 0; i < numElements; ++i) + { + if (output_vector[i] != results[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl; + std::cout << "Done\n"; + + return 0; +}