From 749b9dfaac19588721dd6b200a0e06f0be1d405d Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Wed, 1 Apr 2026 13:59:40 -0400 Subject: [PATCH 1/8] Fix header include path --- include/boost/decimal/decimal64_t.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/boost/decimal/decimal64_t.hpp b/include/boost/decimal/decimal64_t.hpp index f74b09cc4..b3a745afc 100644 --- a/include/boost/decimal/decimal64_t.hpp +++ b/include/boost/decimal/decimal64_t.hpp @@ -10,7 +10,7 @@ #include #include #include -#include "detail/int128.hpp" +#include #include #include #include From 3190f03d23e9807167439ada3661e05c85907a91 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Wed, 1 Apr 2026 15:04:39 -0400 Subject: [PATCH 2/8] Add CUDA annotations for basic decimal64_t functionality --- include/boost/decimal/decimal64_t.hpp | 386 +++++++++++----------- include/boost/decimal/detail/div_impl.hpp | 2 +- include/boost/decimal/detail/mul_impl.hpp | 6 +- 3 files changed, 197 insertions(+), 197 deletions(-) diff --git a/include/boost/decimal/decimal64_t.hpp b/include/boost/decimal/decimal64_t.hpp index b3a745afc..d381582a9 100644 --- a/include/boost/decimal/decimal64_t.hpp +++ b/include/boost/decimal/decimal64_t.hpp @@ -128,22 +128,22 @@ BOOST_DECIMAL_EXPORT class decimal64_t final std::uint64_t bits_ {}; // Returns the un-biased (quantum) exponent - constexpr auto unbiased_exponent() const noexcept -> exponent_type; + BOOST_DECIMAL_CUDA_CONSTEXPR auto unbiased_exponent() const noexcept -> exponent_type; // Returns the biased exponent - constexpr auto biased_exponent() const noexcept -> biased_exponent_type; + BOOST_DECIMAL_CUDA_CONSTEXPR auto biased_exponent() const noexcept -> biased_exponent_type; // Allows direct editing of the exp template - constexpr auto edit_exponent(T exp) noexcept + BOOST_DECIMAL_CUDA_CONSTEXPR auto edit_exponent(T exp) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, T, void); // Returns the significand complete with the bits implied from the combination field - constexpr auto full_significand() const noexcept -> significand_type; - constexpr auto isneg() const noexcept -> bool; - constexpr auto edit_sign(bool sign) noexcept -> void; + BOOST_DECIMAL_CUDA_CONSTEXPR auto full_significand() const noexcept -> significand_type; + BOOST_DECIMAL_CUDA_CONSTEXPR auto isneg() const noexcept -> bool; + BOOST_DECIMAL_CUDA_CONSTEXPR auto edit_sign(bool sign) noexcept -> void; - constexpr auto to_components() const noexcept -> detail::decimal64_t_components; + BOOST_DECIMAL_CUDA_CONSTEXPR auto to_components() const noexcept -> detail::decimal64_t_components; // Attempts conversion to integral type: // If this is nan sets errno to EINVAL and returns 0 @@ -160,8 +160,8 @@ BOOST_DECIMAL_EXPORT class decimal64_t final friend constexpr auto to_decimal(Decimal val) noexcept -> TargetType; // Debug bit pattern - friend constexpr auto from_bits(std::uint64_t bits) noexcept -> decimal64_t; - friend constexpr auto to_bits(decimal64_t rhs) noexcept -> std::uint64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto from_bits(std::uint64_t bits) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto to_bits(decimal64_t rhs) noexcept -> std::uint64_t; // Equality template between any integer type and decimal64_t template @@ -183,7 +183,7 @@ BOOST_DECIMAL_EXPORT class decimal64_t final -> std::enable_if_t<(detail::is_decimal_floating_point_v && detail::is_decimal_floating_point_v), bool>; - friend constexpr auto d64_div_impl(decimal64_t lhs, decimal64_t rhs, decimal64_t& q, decimal64_t& r) noexcept -> void; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_div_impl(decimal64_t lhs, decimal64_t rhs, decimal64_t& q, decimal64_t& r) noexcept -> void; template friend constexpr auto ilogb(T d) noexcept @@ -195,7 +195,7 @@ BOOST_DECIMAL_EXPORT class decimal64_t final // Micro-optimization: Nearly every call to isfinite in the basic operators is !isfinite. // We can super easily combine this into a single operation - friend constexpr auto not_finite(decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto not_finite(decimal64_t rhs) noexcept -> bool; template friend constexpr auto equality_impl(DecimalType lhs, DecimalType rhs) noexcept -> bool; @@ -241,7 +241,7 @@ BOOST_DECIMAL_EXPORT class decimal64_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 decimal64_t value) noexcept -> decimal64_t + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto nan_conversion(const decimal64_t value) noexcept -> decimal64_t { constexpr auto convert_nan_mask {detail::d64_snan_mask ^ detail::d64_nan_mask}; @@ -299,33 +299,33 @@ BOOST_DECIMAL_EXPORT class decimal64_t final #if !defined(BOOST_DECIMAL_ALLOW_IMPLICIT_CONVERSIONS) && !defined(BOOST_DECIMAL_ALLOW_IMPLICIT_INTEGER_CONVERSIONS) explicit #endif - constexpr decimal64_t(Integer val) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_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, decimal64_t&); // 3.2.3.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; #ifdef BOOST_DECIMAL_HAS_INT128 - explicit constexpr operator detail::builtin_int128_t() const noexcept; - explicit constexpr operator detail::builtin_uint128_t() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator detail::builtin_int128_t() const noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR operator detail::builtin_uint128_t() const noexcept; #endif // Conversion to another decimal type 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; // 3.2.6 Conversion to a floating-point type explicit BOOST_DECIMAL_CXX20_CONSTEXPR operator float() const noexcept; @@ -354,7 +354,7 @@ BOOST_DECIMAL_EXPORT class decimal64_t final #else template && detail::is_integral_v, bool> = true> #endif - constexpr decimal64_t(T1 coeff, T2 exp, detail::construction_sign_wrapper resultant_sign = construction_sign::positive) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t(T1 coeff, T2 exp, detail::construction_sign_wrapper resultant_sign = construction_sign::positive) noexcept; #ifdef BOOST_DECIMAL_HAS_CONCEPTS template @@ -368,9 +368,9 @@ BOOST_DECIMAL_EXPORT class decimal64_t final #else template && detail::is_integral_v, bool> = true> #endif - constexpr decimal64_t(T1 coeff, T2 exp) noexcept; + BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t(T1 coeff, T2 exp) noexcept; - explicit constexpr decimal64_t(bool value) noexcept; + explicit BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t(bool value) noexcept; #if !defined(BOOST_DECIMAL_DISABLE_CLIB) @@ -385,212 +385,212 @@ BOOST_DECIMAL_EXPORT class decimal64_t final #endif // cmath functions that are easier as friends - friend constexpr auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; - friend constexpr auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; - friend constexpr auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; - friend constexpr auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; - friend constexpr auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; - friend constexpr auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (decimal64_t rhs) noexcept -> bool; // 3.2.7 unary arithmetic operators: - friend constexpr auto operator+(decimal64_t rhs) noexcept -> decimal64_t; - friend constexpr auto operator-(decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal64_t rhs) noexcept -> decimal64_t; // 3.2.8 Binary arithmetic operators - friend constexpr auto operator+(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; template - friend constexpr auto operator+(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); template - friend constexpr auto operator+(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); - friend constexpr auto operator-(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; template - friend constexpr auto operator-(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); template - friend constexpr auto operator-(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); - friend constexpr auto operator*(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; template - friend constexpr auto operator*(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); template - friend constexpr auto operator*(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); - friend constexpr auto operator/(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; template - friend constexpr auto operator/(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); template - friend constexpr auto operator/(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t); - friend constexpr auto operator%(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator%(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; // 3.2.3.5 Increment and Decrement - constexpr auto operator++() noexcept -> decimal64_t&; - constexpr auto operator++(int) noexcept -> decimal64_t; // NOLINT : C++14 so constexpr implies const - constexpr auto operator--() noexcept -> decimal64_t&; - constexpr auto operator--(int) noexcept -> decimal64_t; // NOLINT : C++14 so constexpr implies const + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator++() noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator++(int) noexcept -> decimal64_t; // NOLINT : C++14 so constexpr implies const + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator--() noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator--(int) noexcept -> decimal64_t; // NOLINT : C++14 so constexpr implies const // 3.2.3.6 Compound assignment - constexpr auto operator+=(decimal64_t rhs) noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+=(decimal64_t rhs) noexcept -> decimal64_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, decimal64_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, decimal64_t&); - constexpr auto operator-=(decimal64_t rhs) noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-=(decimal64_t rhs) noexcept -> decimal64_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, decimal64_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, decimal64_t&); - constexpr auto operator*=(decimal64_t rhs) noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*=(decimal64_t rhs) noexcept -> decimal64_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, decimal64_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, decimal64_t&); - constexpr auto operator/=(decimal64_t rhs) noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/=(decimal64_t rhs) noexcept -> decimal64_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, decimal64_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, decimal64_t&); - constexpr auto operator%=(decimal64_t rhs) noexcept -> decimal64_t&; + BOOST_DECIMAL_CUDA_CONSTEXPR auto operator%=(decimal64_t rhs) noexcept -> decimal64_t&; // 3.2.9 Comparison operators: // Equality - friend constexpr auto operator==(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; template - friend constexpr auto operator==(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator==(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Inequality - friend constexpr auto operator!=(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; template - friend constexpr auto operator!=(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator!=(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Less - friend constexpr auto operator<(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; template - friend constexpr auto operator<(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator<(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Less equal - friend constexpr auto operator<=(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; template - friend constexpr auto operator<=(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator<=(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Greater - friend constexpr auto operator>(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; template - friend constexpr auto operator>(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator>(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); // Greater equal - friend constexpr auto operator>=(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; template - friend constexpr auto operator>=(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool); template - friend constexpr auto operator>=(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(Integer lhs, decimal64_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<=>(decimal64_t lhs, decimal64_t rhs) noexcept -> std::partial_ordering; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(decimal64_t lhs, decimal64_t rhs) noexcept -> std::partial_ordering; template - friend constexpr auto operator<=>(decimal64_t lhs, Integer rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(decimal64_t lhs, Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering); template - friend constexpr auto operator<=>(Integer lhs, decimal64_t rhs) noexcept + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(Integer lhs, decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering); #endif // 3.6.4 Same Quantum - friend constexpr auto samequantumd64(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto samequantumd64(decimal64_t lhs, decimal64_t rhs) noexcept -> bool; // 3.6.5 Quantum exponent - friend constexpr auto quantexpd64(decimal64_t x) noexcept -> int; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto quantexpd64(decimal64_t x) noexcept -> int; // 3.6.6 Quantize - friend constexpr auto quantized64(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto quantized64(decimal64_t lhs, decimal64_t rhs) noexcept -> decimal64_t; // functions that need to be friends template friend constexpr auto frexp10(T num, int* expptr) noexcept -> typename T::significand_type; - friend constexpr auto copysignd64(decimal64_t mag, decimal64_t sgn) noexcept -> decimal64_t; - friend constexpr auto scalbnd64(decimal64_t num, int exp) noexcept -> decimal64_t; - friend constexpr auto scalblnd64(decimal64_t num, long exp) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto copysignd64(decimal64_t mag, decimal64_t sgn) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto scalbnd64(decimal64_t num, int exp) noexcept -> decimal64_t; + friend BOOST_DECIMAL_CUDA_CONSTEXPR auto scalblnd64(decimal64_t num, long exp) noexcept -> decimal64_t; }; #if defined(__GNUC__) && __GNUC__ >= 8 # pragma GCC diagnostic pop #endif -constexpr auto from_bits(std::uint64_t bits) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto from_bits(std::uint64_t bits) noexcept -> decimal64_t { decimal64_t result; result.bits_ = bits; @@ -616,7 +616,7 @@ template #else template && detail::is_integral_v, bool>> #endif -constexpr decimal64_t::decimal64_t(T1 coeff, T2 exp, const detail::construction_sign_wrapper resultant_sign) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::decimal64_t(T1 coeff, T2 exp, const detail::construction_sign_wrapper resultant_sign) noexcept { const auto is_negative {static_cast(resultant_sign)}; bits_ = is_negative ? detail::d64_sign_mask : UINT64_C(0); @@ -800,9 +800,9 @@ template #else template && detail::is_integral_v, bool>> #endif -constexpr decimal64_t::decimal64_t(const T1 coeff, const T2 exp) noexcept : decimal64_t(detail::make_positive_unsigned(coeff), exp, coeff < 0) {} +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::decimal64_t(const T1 coeff, const T2 exp) noexcept : decimal64_t(detail::make_positive_unsigned(coeff), exp, coeff < 0) {} -constexpr decimal64_t::decimal64_t(const bool value) noexcept : decimal64_t(static_cast(value), 0, false) {} +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::decimal64_t(const bool value) noexcept : decimal64_t(static_cast(value), 0, false) {} #if defined(__GNUC__) && __GNUC__ >= 6 @@ -846,15 +846,15 @@ class numeric_limits_impl64 static constexpr bool tinyness_before = true; // Member functions - static constexpr auto (min) () -> boost::decimal::decimal64_t { return {UINT32_C(1), min_exponent}; } - static constexpr auto (max) () -> boost::decimal::decimal64_t { return {boost::decimal::detail::d64_max_significand_value, max_exponent - digits + 1}; } - static constexpr auto lowest () -> boost::decimal::decimal64_t { return {boost::decimal::detail::d64_max_significand_value, max_exponent - digits + 1, construction_sign::negative}; } - static constexpr auto epsilon () -> boost::decimal::decimal64_t { return {UINT32_C(1), -digits + 1}; } - static constexpr auto round_error () -> boost::decimal::decimal64_t { return epsilon(); } - static constexpr auto infinity () -> boost::decimal::decimal64_t { return boost::decimal::from_bits(boost::decimal::detail::d64_inf_mask); } - static constexpr auto quiet_NaN () -> boost::decimal::decimal64_t { return boost::decimal::from_bits(boost::decimal::detail::d64_nan_mask); } - static constexpr auto signaling_NaN() -> boost::decimal::decimal64_t { return boost::decimal::from_bits(boost::decimal::detail::d64_snan_mask); } - static constexpr auto denorm_min () -> boost::decimal::decimal64_t { return {1, boost::decimal::detail::etiny_v}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto (min) () -> boost::decimal::decimal64_t { return {UINT32_C(1), min_exponent}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto (max) () -> boost::decimal::decimal64_t { return {boost::decimal::detail::d64_max_significand_value, max_exponent - digits + 1}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto lowest () -> boost::decimal::decimal64_t { return {boost::decimal::detail::d64_max_significand_value, max_exponent - digits + 1, construction_sign::negative}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto epsilon () -> boost::decimal::decimal64_t { return {UINT32_C(1), -digits + 1}; } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto round_error () -> boost::decimal::decimal64_t { return epsilon(); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto infinity () -> boost::decimal::decimal64_t { return boost::decimal::from_bits(boost::decimal::detail::d64_inf_mask); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto quiet_NaN () -> boost::decimal::decimal64_t { return boost::decimal::from_bits(boost::decimal::detail::d64_nan_mask); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto signaling_NaN() -> boost::decimal::decimal64_t { return boost::decimal::from_bits(boost::decimal::detail::d64_snan_mask); } + static BOOST_DECIMAL_CUDA_CONSTEXPR auto denorm_min () -> boost::decimal::decimal64_t { return {1, boost::decimal::detail::etiny_v}; } }; #if !defined(__cpp_inline_variables) || __cpp_inline_variables < 201606L @@ -990,10 +990,10 @@ template #else template , bool>> #endif -constexpr decimal64_t::decimal64_t(const Integer val) noexcept : decimal64_t{val, 0} {} +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::decimal64_t(const Integer val) noexcept : decimal64_t{val, 0} {} template -constexpr auto decimal64_t::operator=(const Integer& val) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator=(const Integer& val) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t&) { using ConversionType = std::conditional_t::value, std::int32_t, Integer>; @@ -1001,50 +1001,50 @@ constexpr auto decimal64_t::operator=(const Integer& val) noexcept return *this; } -constexpr decimal64_t::operator bool() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator bool() const noexcept { constexpr decimal64_t zero {0, 0}; return *this != zero; } -constexpr decimal64_t::operator int() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator int() const noexcept { return to_integral(*this); } -constexpr decimal64_t::operator unsigned() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator unsigned() const noexcept { return to_integral(*this); } -constexpr decimal64_t::operator long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator long() const noexcept { return to_integral(*this); } -constexpr decimal64_t::operator unsigned long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator unsigned long() const noexcept { return to_integral(*this); } -constexpr decimal64_t::operator long long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator long long() const noexcept { return to_integral(*this); } -constexpr decimal64_t::operator unsigned long long() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator unsigned long long() const noexcept { return to_integral(*this); } #ifdef BOOST_DECIMAL_HAS_INT128 -constexpr decimal64_t::operator detail::builtin_int128_t() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator detail::builtin_int128_t() const noexcept { return to_integral(*this); } -constexpr decimal64_t::operator detail::builtin_uint128_t() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator detail::builtin_uint128_t() const noexcept { return to_integral(*this); } @@ -1052,13 +1052,13 @@ constexpr decimal64_t::operator detail::builtin_uint128_t() const noexcept #endif template && (detail::decimal_val_v > detail::decimal_val_v), bool>> -constexpr decimal64_t::operator Decimal() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator Decimal() const noexcept { return to_decimal(*this); } template && (detail::decimal_val_v <= detail::decimal_val_v), bool>> -constexpr decimal64_t::operator Decimal() const noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR decimal64_t::operator Decimal() const noexcept { return to_decimal(*this); } @@ -1105,7 +1105,7 @@ constexpr decimal64_t::operator std::bfloat16_t() const noexcept } #endif -constexpr auto decimal64_t::unbiased_exponent() const noexcept -> exponent_type +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::unbiased_exponent() const noexcept -> exponent_type { exponent_type expval {}; @@ -1121,12 +1121,12 @@ constexpr auto decimal64_t::unbiased_exponent() const noexcept -> exponent_type return expval; } -constexpr auto decimal64_t::biased_exponent() const noexcept -> biased_exponent_type +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::biased_exponent() const noexcept -> biased_exponent_type { return static_cast(unbiased_exponent()) - detail::bias_v; } -constexpr auto decimal64_t::full_significand() const noexcept -> significand_type +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::full_significand() const noexcept -> significand_type { significand_type significand {}; @@ -1143,12 +1143,12 @@ constexpr auto decimal64_t::full_significand() const noexcept -> significand_typ return significand; } -constexpr auto decimal64_t::isneg() const noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::isneg() const noexcept -> bool { return static_cast(bits_ & detail::d64_sign_mask); } -constexpr auto decimal64_t::to_components() const noexcept -> detail::decimal64_t_components +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::to_components() const noexcept -> detail::decimal64_t_components { detail::decimal64_t_components components {}; @@ -1175,13 +1175,13 @@ constexpr auto decimal64_t::to_components() const noexcept -> detail::decimal64_ } template -constexpr auto decimal64_t::edit_exponent(const T expval) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::edit_exponent(const T expval) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, T, void) { *this = decimal64_t(this->full_significand(), expval, this->isneg()); } -constexpr auto decimal64_t::edit_sign(const bool sign) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::edit_sign(const bool sign) noexcept -> void { if (sign) { @@ -1193,12 +1193,12 @@ constexpr auto decimal64_t::edit_sign(const bool sign) noexcept -> void } } -constexpr auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto signbit BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool { return rhs.bits_ & detail::d64_sign_mask; } -constexpr auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_ & detail::d64_nan_mask) == detail::d64_nan_mask; @@ -1208,7 +1208,7 @@ constexpr auto isnan BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t #endif } -constexpr auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return ((rhs.bits_ & detail::d64_nan_mask) == detail::d64_inf_mask); @@ -1218,7 +1218,7 @@ constexpr auto isinf BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t #endif } -constexpr auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return (rhs.bits_ & detail::d64_snan_mask) == detail::d64_snan_mask; @@ -1228,7 +1228,7 @@ constexpr auto issignaling BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decim #endif } -constexpr auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH // Check for de-normals @@ -1246,7 +1246,7 @@ constexpr auto isnormal BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal6 #endif } -constexpr auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return ((rhs.bits_ & detail::d64_inf_mask) != detail::d64_inf_mask); @@ -1256,7 +1256,7 @@ constexpr auto isfinite BOOST_DECIMAL_PREVENT_MACRO_SUBSTITUTION (const decimal6 #endif } -BOOST_DECIMAL_FORCE_INLINE constexpr auto not_finite(const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto not_finite(const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH return ((rhs.bits_ & detail::d64_inf_mask) == detail::d64_inf_mask); @@ -1266,45 +1266,45 @@ BOOST_DECIMAL_FORCE_INLINE constexpr auto not_finite(const decimal64_t rhs) noex #endif } -constexpr auto operator==(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { return equality_impl(lhs, rhs); } template -constexpr auto operator==(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const decimal64_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 decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator==(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return mixed_equality_impl(rhs, lhs); } -constexpr auto operator!=(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { return !(lhs == rhs); } template -constexpr auto operator!=(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const decimal64_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 decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator!=(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return !(lhs == rhs); } -constexpr auto operator<(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1329,14 +1329,14 @@ constexpr auto operator<(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator<(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const decimal64_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 decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1349,7 +1349,7 @@ constexpr auto operator<(const Integer lhs, const decimal64_t rhs) noexcept return !less_impl(rhs, lhs) && lhs != rhs; } -constexpr auto operator<=(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH if (isnan(lhs) || isnan(rhs)) @@ -1362,7 +1362,7 @@ constexpr auto operator<=(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator<=(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1376,7 +1376,7 @@ constexpr auto operator<=(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator<=(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1389,26 +1389,26 @@ constexpr auto operator<=(const Integer lhs, const decimal64_t rhs) noexcept return !(rhs < lhs); } -constexpr auto operator>(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { return rhs < lhs; } template -constexpr auto operator>(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const decimal64_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 decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { return rhs < lhs; } -constexpr auto operator>=(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH if (isnan(lhs) || isnan(rhs)) @@ -1421,7 +1421,7 @@ constexpr auto operator>=(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator>=(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1435,7 +1435,7 @@ constexpr auto operator>=(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator>=(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator>=(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, bool) { #ifndef BOOST_DECIMAL_FAST_MATH @@ -1450,7 +1450,7 @@ constexpr auto operator>=(const Integer lhs, const decimal64_t rhs) noexcept #ifdef BOOST_DECIMAL_HAS_SPACESHIP_OPERATOR -constexpr auto operator<=>(const decimal64_t lhs, const decimal64_t rhs) noexcept -> std::partial_ordering +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const decimal64_t lhs, const decimal64_t rhs) noexcept -> std::partial_ordering { if (lhs < rhs) { @@ -1469,7 +1469,7 @@ constexpr auto operator<=>(const decimal64_t lhs, const decimal64_t rhs) noexcep } template -constexpr auto operator<=>(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering) { if (lhs < rhs) @@ -1489,7 +1489,7 @@ constexpr auto operator<=>(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator<=>(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator<=>(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, std::partial_ordering) { if (lhs < rhs) @@ -1510,18 +1510,18 @@ constexpr auto operator<=>(const Integer lhs, const decimal64_t rhs) noexcept #endif -constexpr auto operator+(const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal64_t rhs) noexcept -> decimal64_t { return rhs; } -constexpr auto operator-(decimal64_t rhs) noexcept-> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(decimal64_t rhs) noexcept-> decimal64_t { rhs.bits_ ^= detail::d64_sign_mask; return rhs; } -constexpr auto d64_div_impl(const decimal64_t lhs, const decimal64_t rhs, decimal64_t& q, decimal64_t& r) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_div_impl(const decimal64_t lhs, const decimal64_t rhs, decimal64_t& q, decimal64_t& r) noexcept -> void { const bool sign {lhs.isneg() != rhs.isneg()}; @@ -1632,7 +1632,7 @@ constexpr auto d64_div_impl(const decimal64_t lhs, const decimal64_t rhs, decima q = detail::d64_generic_div_impl(lhs_components, rhs.to_components(), sign); } -constexpr auto operator+(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1655,7 +1655,7 @@ constexpr auto operator+(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator+(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { using promoted_significand_type = detail::promote_significand_t; @@ -1684,14 +1684,14 @@ constexpr auto operator+(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator+(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator+(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { return rhs + lhs; } // NOLINTNEXTLINE: If subtraction is actually addition than use operator+ and vice versa -constexpr auto operator-(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1719,7 +1719,7 @@ constexpr auto operator-(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator-(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { using promoted_significand_type = detail::promote_significand_t; @@ -1748,7 +1748,7 @@ constexpr auto operator-(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator-(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator-(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { using promoted_significand_type = detail::promote_significand_t; @@ -1781,7 +1781,7 @@ constexpr auto operator-(const Integer lhs, const decimal64_t rhs) noexcept ); } -constexpr auto operator*(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(lhs) || not_finite(rhs)) @@ -1818,7 +1818,7 @@ constexpr auto operator*(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator*(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { using promoted_significand_type = detail::promote_significand_t; @@ -1854,13 +1854,13 @@ constexpr auto operator*(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator*(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator*(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { return rhs * lhs; } -constexpr auto operator/(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t { decimal64_t q {}; decimal64_t r {}; @@ -1870,7 +1870,7 @@ constexpr auto operator/(const decimal64_t lhs, const decimal64_t rhs) noexcept } template -constexpr auto operator/(const decimal64_t lhs, const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const decimal64_t lhs, const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { using sig_type = decimal64_t::significand_type; @@ -1918,7 +1918,7 @@ constexpr auto operator/(const decimal64_t lhs, const Integer rhs) noexcept } template -constexpr auto operator/(const Integer lhs, const decimal64_t rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator/(const Integer lhs, const decimal64_t rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t) { using sig_type = decimal64_t::significand_type; @@ -1960,7 +1960,7 @@ constexpr auto operator/(const Integer lhs, const decimal64_t rhs) noexcept return detail::d64_generic_div_impl(lhs_components, rhs_components, sign); } -constexpr auto operator%(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto operator%(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t { decimal64_t q {}; decimal64_t r {}; @@ -2006,42 +2006,42 @@ constexpr auto operator%(const decimal64_t lhs, const decimal64_t rhs) noexcept return r; } -constexpr auto decimal64_t::operator++() noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator++() noexcept -> decimal64_t& { constexpr decimal64_t one{1, 0}; *this = *this + one; return *this; } -constexpr auto decimal64_t::operator++(int) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator++(int) noexcept -> decimal64_t { const auto temp {*this}; ++(*this); return temp; } -constexpr auto decimal64_t::operator--() noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator--() noexcept -> decimal64_t& { constexpr decimal64_t one{1, 0}; *this = *this - one; return *this; } -constexpr auto decimal64_t::operator--(int) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator--(int) noexcept -> decimal64_t { const auto temp {*this}; --(*this); return temp; } -constexpr auto decimal64_t::operator+=(const decimal64_t rhs) noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator+=(const decimal64_t rhs) noexcept -> decimal64_t& { *this = *this + rhs; return *this; } template -constexpr auto decimal64_t::operator+=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator+=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t&) { *this = *this + rhs; @@ -2049,21 +2049,21 @@ constexpr auto decimal64_t::operator+=(const Integer rhs) noexcept } template -constexpr auto decimal64_t::operator+=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator+=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal64_t&) { *this = *this + rhs; return *this; } -constexpr auto decimal64_t::operator-=(const decimal64_t rhs) noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator-=(const decimal64_t rhs) noexcept -> decimal64_t& { *this = *this - rhs; return *this; } template -constexpr auto decimal64_t::operator-=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator-=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t&) { *this = *this - rhs; @@ -2071,21 +2071,21 @@ constexpr auto decimal64_t::operator-=(const Integer rhs) noexcept } template -constexpr auto decimal64_t::operator-=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator-=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal64_t&) { *this = *this - rhs; return *this; } -constexpr auto decimal64_t::operator*=(const decimal64_t rhs) noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator*=(const decimal64_t rhs) noexcept -> decimal64_t& { *this = *this * rhs; return *this; } template -constexpr auto decimal64_t::operator*=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator*=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t&) { *this = *this * rhs; @@ -2093,21 +2093,21 @@ constexpr auto decimal64_t::operator*=(const Integer rhs) noexcept } template -constexpr auto decimal64_t::operator*=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator*=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal64_t&) { *this = *this * rhs; return *this; } -constexpr auto decimal64_t::operator/=(const decimal64_t rhs) noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator/=(const decimal64_t rhs) noexcept -> decimal64_t& { *this = *this / rhs; return *this; } template -constexpr auto decimal64_t::operator/=(const Integer rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator/=(const Integer rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_integral_v, Integer, decimal64_t&) { *this = *this / rhs; @@ -2115,14 +2115,14 @@ constexpr auto decimal64_t::operator/=(const Integer rhs) noexcept } template -constexpr auto decimal64_t::operator/=(const Decimal rhs) noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator/=(const Decimal rhs) noexcept BOOST_DECIMAL_REQUIRES_RETURN(detail::is_decimal_floating_point_v, Decimal, decimal64_t&) { *this = *this / rhs; return *this; } -constexpr auto decimal64_t::operator%=(const decimal64_t rhs) noexcept -> decimal64_t& +BOOST_DECIMAL_CUDA_CONSTEXPR auto decimal64_t::operator%=(const decimal64_t rhs) noexcept -> decimal64_t& { *this = *this % rhs; return *this; @@ -2133,7 +2133,7 @@ constexpr auto decimal64_t::operator%=(const decimal64_t rhs) noexcept -> decima // 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 samequantumd64(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto samequantumd64(const decimal64_t lhs, const decimal64_t rhs) noexcept -> bool { #ifndef BOOST_DECIMAL_FAST_MATH const auto lhs_fp {fpclassify(lhs)}; @@ -2155,7 +2155,7 @@ constexpr auto samequantumd64(const decimal64_t lhs, const decimal64_t rhs) noex // 3.6.5 // Effects: if x is finite, returns its quantum exponent. // Otherwise, a domain error occurs and INT_MIN is returned. -constexpr auto quantexpd64(const decimal64_t x) noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto quantexpd64(const decimal64_t x) noexcept -> int { #ifndef BOOST_DECIMAL_FAST_MATH if (not_finite(x)) @@ -2178,7 +2178,7 @@ constexpr auto quantexpd64(const decimal64_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 quantized64(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto quantized64(const decimal64_t lhs, const decimal64_t rhs) noexcept -> decimal64_t { #ifndef BOOST_DECIMAL_FAST_MATH // Return the correct type of nan @@ -2205,7 +2205,7 @@ constexpr auto quantized64(const decimal64_t lhs, const decimal64_t rhs) noexcep return {lhs.full_significand(), rhs.biased_exponent(), lhs.isneg()}; } -constexpr auto scalblnd64(decimal64_t num, const long exp) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto scalblnd64(decimal64_t num, const long exp) noexcept -> decimal64_t { #ifndef BOOST_DECIMAL_FAST_MATH constexpr decimal64_t zero {0, 0}; @@ -2221,12 +2221,12 @@ constexpr auto scalblnd64(decimal64_t num, const long exp) noexcept -> decimal64 return num; } -constexpr auto scalbnd64(const decimal64_t num, const int expval) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto scalbnd64(const decimal64_t num, const int expval) noexcept -> decimal64_t { return scalblnd64(num, static_cast(expval)); } -constexpr auto copysignd64(decimal64_t mag, const decimal64_t sgn) noexcept -> decimal64_t +BOOST_DECIMAL_CUDA_CONSTEXPR auto copysignd64(decimal64_t mag, const decimal64_t sgn) noexcept -> decimal64_t { mag.edit_sign(sgn.isneg()); return mag; diff --git a/include/boost/decimal/detail/div_impl.hpp b/include/boost/decimal/detail/div_impl.hpp index 3c21522e8..877687d00 100644 --- a/include/boost/decimal/detail/div_impl.hpp +++ b/include/boost/decimal/detail/div_impl.hpp @@ -49,7 +49,7 @@ BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto generic_div_impl(co } template -BOOST_DECIMAL_FORCE_INLINE constexpr auto d64_generic_div_impl(const T& lhs, const T& rhs, const bool sign) noexcept -> DecimalType +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_generic_div_impl(const T& lhs, const T& rhs, const bool sign) noexcept -> DecimalType { using unsigned_int128_type = boost::int128::uint128_t; diff --git a/include/boost/decimal/detail/mul_impl.hpp b/include/boost/decimal/detail/mul_impl.hpp index 4324230f5..9763d4b57 100644 --- a/include/boost/decimal/detail/mul_impl.hpp +++ b/include/boost/decimal/detail/mul_impl.hpp @@ -64,7 +64,7 @@ BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto mul_impl(T lhs_sig, // In the fast case we are better served doing our 128-bit division here since we are at a know starting point template -constexpr auto d64_mul_impl(const T& lhs, const T& rhs) noexcept -> ReturnType +BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_mul_impl(const T& lhs, const T& rhs) noexcept -> ReturnType { using unsigned_int128_type = boost::int128::uint128_t; @@ -85,7 +85,7 @@ constexpr auto d64_mul_impl(const T& lhs, const T& rhs) noexcept -> ReturnType } template -BOOST_DECIMAL_FORCE_INLINE constexpr auto d64_mul_impl(T lhs_sig, U lhs_exp, bool lhs_sign, +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_mul_impl(T lhs_sig, U lhs_exp, bool lhs_sign, T rhs_sig, U rhs_exp, bool rhs_sign) noexcept -> std::enable_if_t, ReturnType> { @@ -108,7 +108,7 @@ BOOST_DECIMAL_FORCE_INLINE constexpr auto d64_mul_impl(T lhs_sig, U lhs_exp, boo } template -BOOST_DECIMAL_FORCE_INLINE constexpr auto d64_mul_impl(T lhs_sig, U lhs_exp, bool lhs_sign, +BOOST_DECIMAL_FORCE_INLINE BOOST_DECIMAL_CUDA_CONSTEXPR auto d64_mul_impl(T lhs_sig, U lhs_exp, bool lhs_sign, T rhs_sig, U rhs_exp, bool rhs_sign) noexcept -> std::enable_if_t, ReturnType> { From e2ee3c58a945ab0666ef6cf4db6b79264be690c6 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Wed, 1 Apr 2026 15:04:56 -0400 Subject: [PATCH 3/8] Add testing of basic decimal64_t functionality on device --- test/cuda_jamfile | 6 ++ test/test_cuda_decimal64_t_add.cu | 76 ++++++++++++++++++++++++ test/test_cuda_decimal64_t_construct.cu | 76 ++++++++++++++++++++++++ test/test_cuda_decimal64_t_div.cu | 79 +++++++++++++++++++++++++ test/test_cuda_decimal64_t_mul.cu | 78 ++++++++++++++++++++++++ test/test_cuda_decimal64_t_sub.cu | 78 ++++++++++++++++++++++++ 6 files changed, 393 insertions(+) create mode 100644 test/test_cuda_decimal64_t_add.cu create mode 100644 test/test_cuda_decimal64_t_construct.cu create mode 100644 test/test_cuda_decimal64_t_div.cu create mode 100644 test/test_cuda_decimal64_t_mul.cu create mode 100644 test/test_cuda_decimal64_t_sub.cu diff --git a/test/cuda_jamfile b/test/cuda_jamfile index 8ad9ca9a0..c8122f67a 100644 --- a/test/cuda_jamfile +++ b/test/cuda_jamfile @@ -9,3 +9,9 @@ run test_cuda_decimal32_t_add.cu ; run test_cuda_decimal32_t_sub.cu ; run test_cuda_decimal32_t_mul.cu ; run test_cuda_decimal32_t_div.cu ; + +run test_cuda_decimal64_t_construct.cu ; +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 ; diff --git a/test/test_cuda_decimal64_t_add.cu b/test/test_cuda_decimal64_t_add.cu new file mode 100644 index 000000000..d61ca364a --- /dev/null +++ b/test/test_cuda_decimal64_t_add.cu @@ -0,0 +1,76 @@ +// 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::decimal64_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; + + 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_decimal64_t_construct.cu b/test/test_cuda_decimal64_t_construct.cu new file mode 100644 index 000000000..ac55293da --- /dev/null +++ b/test/test_cuda_decimal64_t_construct.cu @@ -0,0 +1,76 @@ +// 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::decimal64_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; + + 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_decimal64_t_div.cu b/test/test_cuda_decimal64_t_div.cu new file mode 100644 index 000000000..b052b05b0 --- /dev/null +++ b/test/test_cuda_decimal64_t_div.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::decimal64_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; + + 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_decimal64_t_mul.cu b/test/test_cuda_decimal64_t_mul.cu new file mode 100644 index 000000000..d0fd65f67 --- /dev/null +++ b/test/test_cuda_decimal64_t_mul.cu @@ -0,0 +1,78 @@ +// 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::decimal64_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; + + 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_decimal64_t_sub.cu b/test/test_cuda_decimal64_t_sub.cu new file mode 100644 index 000000000..b8ca86bd8 --- /dev/null +++ b/test/test_cuda_decimal64_t_sub.cu @@ -0,0 +1,78 @@ +// 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::decimal64_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; + + 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; +} From 1d73676a304d9e1a246b3b4fdb03c81f242caa92 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Wed, 1 Apr 2026 15:13:37 -0400 Subject: [PATCH 4/8] Attempt fix illegal memory access error --- include/boost/decimal/detail/attributes.hpp | 24 ++++++++++----------- include/boost/decimal/detail/utilities.hpp | 2 +- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/include/boost/decimal/detail/attributes.hpp b/include/boost/decimal/detail/attributes.hpp index 4f979905a..c7db31f42 100644 --- a/include/boost/decimal/detail/attributes.hpp +++ b/include/boost/decimal/detail/attributes.hpp @@ -25,83 +25,83 @@ namespace detail { namespace impl { template -constexpr auto storage_width_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto storage_width_v() noexcept -> int { return decimal_val_v < 64 ? 32 : decimal_val_v < 128 ? 64 : 128; } template -constexpr auto precision_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto precision_v() noexcept -> int { return decimal_val_v < 64 ? 7 : decimal_val_v < 128 ? 16 : 34; } template -constexpr auto bias_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto bias_v() noexcept -> int { return decimal_val_v < 64 ? 101 : decimal_val_v < 128 ? 398 : 6176; } template -constexpr auto max_biased_exp_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto max_biased_exp_v() noexcept -> int { return decimal_val_v < 64 ? 191 : decimal_val_v < 128 ? 767 : 12287; } template -constexpr auto emax_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto emax_v() noexcept -> int { return decimal_val_v < 64 ? 96 : decimal_val_v < 128 ? 384 : 6144; } template -constexpr auto emin_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto emin_v() noexcept -> int { return decimal_val_v < 64 ? -95 : decimal_val_v < 128 ? -383 : -6143; } template -constexpr auto combination_field_width_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto combination_field_width_v() noexcept -> int { return decimal_val_v < 64 ? 11 : decimal_val_v < 128 ? 13 : 17; } template -constexpr auto trailing_significand_field_width_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto trailing_significand_field_width_v() noexcept -> int { return decimal_val_v < 64 ? 20 : decimal_val_v < 128 ? 50 : 110; } template < 128, bool> = true> -constexpr auto max_significand_v() noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto max_significand_v() noexcept { return decimal_val_v < 64 ? 9'999'999 : 9'999'999'999'999'999; } template >= 128, bool> = true> -constexpr auto max_significand_v() noexcept +BOOST_DECIMAL_CUDA_CONSTEXPR auto max_significand_v() noexcept { // 34x 9s return BOOST_DECIMAL_DETAIL_INT128_UINT128_C(9999999999999999999999999999999999); } template -constexpr auto max_string_length_v() noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto max_string_length_v() noexcept -> int { return decimal_val_v < 64 ? 15 : decimal_val_v < 128 ? 25 : 41; } template -constexpr auto is_fast_type_v() noexcept -> bool +BOOST_DECIMAL_CUDA_CONSTEXPR auto is_fast_type_v() noexcept -> bool { // The fast types all assign 1 additional bit over the regular types return decimal_val_v % 2 == 1; diff --git a/include/boost/decimal/detail/utilities.hpp b/include/boost/decimal/detail/utilities.hpp index 108217cf7..158dc348b 100644 --- a/include/boost/decimal/detail/utilities.hpp +++ b/include/boost/decimal/detail/utilities.hpp @@ -16,7 +16,7 @@ namespace decimal { namespace detail { template -constexpr auto swap(T& x, T& y) noexcept -> void +BOOST_DECIMAL_CUDA_CONSTEXPR auto swap(T& x, T& y) noexcept -> void { const T temp {x}; x = y; From bfa7219cfb6bbc36e51fd0f26cb043f9e0373ed6 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Wed, 1 Apr 2026 15:34:17 -0400 Subject: [PATCH 5/8] Avoid usage of builtin 128-bit integers on device --- include/boost/decimal/detail/fenv_rounding.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/boost/decimal/detail/fenv_rounding.hpp b/include/boost/decimal/detail/fenv_rounding.hpp index 000bdb6ce..9d298b17c 100644 --- a/include/boost/decimal/detail/fenv_rounding.hpp +++ b/include/boost/decimal/detail/fenv_rounding.hpp @@ -61,7 +61,7 @@ BOOST_DECIMAL_CUDA_CONSTEXPR auto divmod(T dividend, T divisor) noexcept -> divm return {q, r}; } -#ifdef BOOST_DECIMAL_DETAIL_INT128_HAS_INT128 +#if defined(BOOST_DECIMAL_DETAIL_INT128_HAS_INT128) && !(defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) constexpr auto divmod(const int128::uint128_t dividend, const int128::uint128_t divisor) -> divmod_result { From a651a675d32712d37bf153cf45082d722beb044f Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Wed, 1 Apr 2026 15:38:01 -0400 Subject: [PATCH 6/8] Remove access to global for rounding --- include/boost/decimal/detail/fenv_rounding.hpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/include/boost/decimal/detail/fenv_rounding.hpp b/include/boost/decimal/detail/fenv_rounding.hpp index 9d298b17c..cbb4bb237 100644 --- a/include/boost/decimal/detail/fenv_rounding.hpp +++ b/include/boost/decimal/detail/fenv_rounding.hpp @@ -101,7 +101,7 @@ BOOST_DECIMAL_CUDA_CONSTEXPR auto divmod10(const int128::uint128_t lhs) noexcept } template -BOOST_DECIMAL_CUDA_CONSTEXPR auto fenv_round_impl(T& val, const bool is_neg, const bool sticky, const rounding_mode round = _boost_decimal_global_rounding_mode) noexcept -> int +BOOST_DECIMAL_CUDA_CONSTEXPR auto fenv_round_impl(T& val, const bool is_neg, const bool sticky, const rounding_mode round) noexcept -> int { using significand_type = std::conditional_t >= 128, int128::uint128_t, std::int64_t>; @@ -166,7 +166,8 @@ BOOST_DECIMAL_CUDA_CONSTEXPR auto fenv_round_impl(T& val, const bool is_neg, con template , bool> = true> BOOST_DECIMAL_CUDA_CONSTEXPR auto fenv_round(T& val, bool is_neg = false, bool sticky = false) noexcept -> int { - return impl::fenv_round_impl(val, is_neg, sticky); + constexpr auto round {_boost_decimal_global_rounding_mode}; + return impl::fenv_round_impl(val, is_neg, sticky, round); } #else @@ -176,7 +177,8 @@ BOOST_DECIMAL_CUDA_CONSTEXPR auto fenv_round(T& val, bool is_neg = false, bool s { if (BOOST_DECIMAL_IS_CONSTANT_EVALUATED(coeff)) { - return impl::fenv_round_impl(val, is_neg, sticky); + constexpr auto round {_boost_decimal_global_rounding_mode}; + return impl::fenv_round_impl(val, is_neg, sticky, round); } else { From 2d41fe487dcb72c17ba69b50061b78800e297578 Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 2 Apr 2026 13:56:02 -0400 Subject: [PATCH 7/8] Match macro configurations --- include/boost/decimal/detail/config.hpp | 14 +++++++------- include/boost/decimal/detail/int128.hpp | 11 ++++++----- 2 files changed, 13 insertions(+), 12 deletions(-) diff --git a/include/boost/decimal/detail/config.hpp b/include/boost/decimal/detail/config.hpp index dcf11862d..9ec5ab623 100644 --- a/include/boost/decimal/detail/config.hpp +++ b/include/boost/decimal/detail/config.hpp @@ -13,8 +13,8 @@ # define BOOST_DECIMAL_DEC_EVAL_METHOD 0 #endif -// Fundamental NVCC options -#if defined(__NVCC__) && defined(BOOST_DECIMAL_ENABLE_CUDA) +// Fundamental CUDACC options +#if defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA) # define BOOST_DECIMAL_HOST_DEVICE __host__ __device__ # define BOOST_DECIMAL_CUDA_CONSTEXPR __host__ __device__ constexpr # define BOOST_DECIMAL_DETAIL_INT128_ENABLE_CUDA @@ -97,7 +97,7 @@ # endif #endif -#if defined(__NVCC__) && defined(BOOST_DECIMAL_ENABLE_CUDA) +#if defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA) // Include intrinsics if available #if defined(_MSC_VER) @@ -147,7 +147,7 @@ #endif // Use 128-bit integers and suppress warnings for using extensions -#if (defined(BOOST_HAS_INT128) || (defined(__SIZEOF_INT128__) && !defined(_MSC_VER))) && !(defined(__NVCC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) +#if (defined(BOOST_HAS_INT128) || (defined(__SIZEOF_INT128__) && !defined(_MSC_VER))) && !(defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) namespace boost { namespace decimal { namespace detail { @@ -170,7 +170,7 @@ typedef unsigned __int128 builtin_uint128_t; #endif // 128-bit floats -#if (defined(BOOST_HAS_FLOAT128) || defined(__SIZEOF_FLOAT128__)) && !(defined(__NVCC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) +#if (defined(BOOST_HAS_FLOAT128) || defined(__SIZEOF_FLOAT128__)) && !(defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) # define BOOST_DECIMAL_HAS_FLOAT128 #endif @@ -283,9 +283,9 @@ typedef unsigned __int128 builtin_uint128_t; # define BOOST_DECIMAL_HAS_BUILTIN_IS_CONSTANT_EVALUATED #endif -#if defined(BOOST_DECIMAL_HAS_IS_CONSTANT_EVALUATED) && !(defined(__NVCC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) +#if defined(BOOST_DECIMAL_HAS_IS_CONSTANT_EVALUATED) && !(defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) # define BOOST_DECIMAL_IS_CONSTANT_EVALUATED(x) std::is_constant_evaluated() -#elif defined(BOOST_DECIMAL_HAS_BUILTIN_IS_CONSTANT_EVALUATED) && !(defined(__NVCC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) +#elif defined(BOOST_DECIMAL_HAS_BUILTIN_IS_CONSTANT_EVALUATED) && !(defined(__CUDACC__) && defined(BOOST_DECIMAL_ENABLE_CUDA)) # define BOOST_DECIMAL_IS_CONSTANT_EVALUATED(x) __builtin_is_constant_evaluated() #else # define BOOST_DECIMAL_IS_CONSTANT_EVALUATED(x) false diff --git a/include/boost/decimal/detail/int128.hpp b/include/boost/decimal/detail/int128.hpp index 5d3fd7ea6..690a0a4bc 100644 --- a/include/boost/decimal/detail/int128.hpp +++ b/include/boost/decimal/detail/int128.hpp @@ -5,10 +5,11 @@ #ifndef BOOST_DECIMAL_DETAIL_INT128_HPP #define BOOST_DECIMAL_DETAIL_INT128_HPP -#include "int128/int128.hpp" -#include "int128/bit.hpp" -#include "int128/iostream.hpp" -#include "int128/literals.hpp" -#include "int128/numeric.hpp" +#include +#include +#include +#include +#include +#include #endif // BOOST_DECIMAL_DETAIL_INT128_HPP From 0bcb79b97de2d6cfd93572aabb8124c7c1cfda9e Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Thu, 2 Apr 2026 14:19:48 -0400 Subject: [PATCH 8/8] Increase the size of the stack from 1024 to 2048 --- test/test_cuda_decimal64_t_add.cu | 1 + test/test_cuda_decimal64_t_div.cu | 1 + test/test_cuda_decimal64_t_mul.cu | 1 + test/test_cuda_decimal64_t_sub.cu | 1 + 4 files changed, 4 insertions(+) diff --git a/test/test_cuda_decimal64_t_add.cu b/test/test_cuda_decimal64_t_add.cu index d61ca364a..71f0fbf59 100644 --- a/test/test_cuda_decimal64_t_add.cu +++ b/test/test_cuda_decimal64_t_add.cu @@ -46,6 +46,7 @@ int main(void) watch w; + cudaDeviceSetLimit(cudaLimitStackSize, 2048); cuda_test<<>>(input_vector.get(), output_vector.get(), numElements); cudaDeviceSynchronize(); diff --git a/test/test_cuda_decimal64_t_div.cu b/test/test_cuda_decimal64_t_div.cu index b052b05b0..c43809690 100644 --- a/test/test_cuda_decimal64_t_div.cu +++ b/test/test_cuda_decimal64_t_div.cu @@ -49,6 +49,7 @@ int main(void) watch w; + cudaDeviceSetLimit(cudaLimitStackSize, 2048); cuda_test<<>>(input_vector.get(), input_vector2.get(), output_vector.get(), numElements); cudaDeviceSynchronize(); diff --git a/test/test_cuda_decimal64_t_mul.cu b/test/test_cuda_decimal64_t_mul.cu index d0fd65f67..ff45e6b79 100644 --- a/test/test_cuda_decimal64_t_mul.cu +++ b/test/test_cuda_decimal64_t_mul.cu @@ -48,6 +48,7 @@ int main(void) watch w; + cudaDeviceSetLimit(cudaLimitStackSize, 2048); cuda_test<<>>(input_vector.get(), input_vector2.get(), output_vector.get(), numElements); cudaDeviceSynchronize(); diff --git a/test/test_cuda_decimal64_t_sub.cu b/test/test_cuda_decimal64_t_sub.cu index b8ca86bd8..a98e5f92d 100644 --- a/test/test_cuda_decimal64_t_sub.cu +++ b/test/test_cuda_decimal64_t_sub.cu @@ -48,6 +48,7 @@ int main(void) watch w; + cudaDeviceSetLimit(cudaLimitStackSize, 2048); cuda_test<<>>(input_vector.get(), input_vector2.get(), output_vector.get(), numElements); cudaDeviceSynchronize();