diff --git a/include/boost/numeric/interval/arith.hpp b/include/boost/numeric/interval/arith.hpp index 0d598be..299e4c2 100644 --- a/include/boost/numeric/interval/arith.hpp +++ b/include/boost/numeric/interval/arith.hpp @@ -26,13 +26,13 @@ namespace numeric { */ template inline -const interval& operator+(const interval& x) +BOOST_GPU_ENABLED const interval& operator+(const interval& x) { return x; } template inline -interval operator-(const interval& x) +BOOST_GPU_ENABLED interval operator-(const interval& x) { if (interval_lib::detail::test_input(x)) return interval::empty(); @@ -40,7 +40,7 @@ interval operator-(const interval& x) } template inline -interval& interval::operator+=(const interval& r) +BOOST_GPU_ENABLED interval& interval::operator+=(const interval& r) { if (interval_lib::detail::test_input(*this, r)) set_empty(); @@ -52,7 +52,7 @@ interval& interval::operator+=(const interval inline -interval& interval::operator+=(const T& r) +BOOST_GPU_ENABLED interval& interval::operator+=(const T& r) { if (interval_lib::detail::test_input(*this, r)) set_empty(); @@ -64,7 +64,7 @@ interval& interval::operator+=(const T& r) } template inline -interval& interval::operator-=(const interval& r) +BOOST_GPU_ENABLED interval& interval::operator-=(const interval& r) { if (interval_lib::detail::test_input(*this, r)) set_empty(); @@ -76,7 +76,7 @@ interval& interval::operator-=(const interval inline -interval& interval::operator-=(const T& r) +BOOST_GPU_ENABLED interval& interval::operator-=(const T& r) { if (interval_lib::detail::test_input(*this, r)) set_empty(); @@ -88,32 +88,32 @@ interval& interval::operator-=(const T& r) } template inline -interval& interval::operator*=(const interval& r) +BOOST_GPU_ENABLED interval& interval::operator*=(const interval& r) { return *this = *this * r; } template inline -interval& interval::operator*=(const T& r) +BOOST_GPU_ENABLED interval& interval::operator*=(const T& r) { return *this = r * *this; } template inline -interval& interval::operator/=(const interval& r) +BOOST_GPU_ENABLED interval& interval::operator/=(const interval& r) { return *this = *this / r; } template inline -interval& interval::operator/=(const T& r) +BOOST_GPU_ENABLED interval& interval::operator/=(const T& r) { return *this = *this / r; } template inline -interval operator+(const interval& x, - const interval& y) +BOOST_GPU_ENABLED interval operator+(const interval& x, + const interval& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -123,7 +123,7 @@ interval operator+(const interval& x, } template inline -interval operator+(const T& x, const interval& y) +BOOST_GPU_ENABLED interval operator+(const T& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -133,11 +133,19 @@ interval operator+(const T& x, const interval& y) } template inline -interval operator+(const interval& x, const T& y) +BOOST_GPU_ENABLED interval operator+(const interval& x, const T& y) { return y + x; } +template inline +BOOST_GPU_ENABLED interval operator+(const T2& x, const interval& y) +{ return static_cast(x) + y; } + +template inline +BOOST_GPU_ENABLED interval operator+(const interval& x, const T2& y) +{ return x + static_cast(y); } + template inline -interval operator-(const interval& x, +BOOST_GPU_ENABLED interval operator-(const interval& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) @@ -147,8 +155,8 @@ interval operator-(const interval& x, rnd.sub_up (x.upper(), y.lower()), true); } -template inline -interval operator-(const T& x, const interval& y) + template inline + BOOST_GPU_ENABLED interval operator-(const T& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -158,7 +166,7 @@ interval operator-(const T& x, const interval& y) } template inline -interval operator-(const interval& x, const T& y) +BOOST_GPU_ENABLED interval operator-(const interval& x, const T& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -167,9 +175,17 @@ interval operator-(const interval& x, const T& y) rnd.sub_up (x.upper(), y), true); } +template inline +BOOST_GPU_ENABLED interval operator-(const T2& x, const interval& y) +{ return static_cast(x) - y; } + +template inline +BOOST_GPU_ENABLED interval operator-(const interval& x, const T2& y) +{ return x - static_cast(y); } + template inline -interval operator*(const interval& x, - const interval& y) +BOOST_GPU_ENABLED interval operator*(const interval& x, + const interval& y) { BOOST_USING_STD_MIN(); BOOST_USING_STD_MAX(); @@ -223,7 +239,7 @@ interval operator*(const interval& x, } template inline -interval operator*(const T& x, const interval& y) +BOOST_GPU_ENABLED interval operator*(const T& x, const interval& y) { typedef interval I; if (interval_lib::detail::test_input(x, y)) @@ -241,11 +257,19 @@ interval operator*(const T& x, const interval& y) } template inline -interval operator*(const interval& x, const T& y) +BOOST_GPU_ENABLED interval operator*(const interval& x, const T& y) { return y * x; } +template inline +BOOST_GPU_ENABLED interval operator*(const T2& x, const interval& y) +{ return static_cast(x) * y; } + +template inline +BOOST_GPU_ENABLED interval operator*(const interval& x, const T2& y) +{ return x * static_cast(y); } + template inline -interval operator/(const interval& x, +BOOST_GPU_ENABLED interval operator/(const interval& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) @@ -266,7 +290,7 @@ interval operator/(const interval& x, } template inline -interval operator/(const T& x, const interval& y) +BOOST_GPU_ENABLED interval operator/(const T& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -286,7 +310,7 @@ interval operator/(const T& x, const interval& y) } template inline -interval operator/(const interval& x, const T& y) +BOOST_GPU_ENABLED interval operator/(const interval& x, const T& y) { if (interval_lib::detail::test_input(x, y) || interval_lib::user::is_zero(y)) return interval::empty(); @@ -299,6 +323,10 @@ interval operator/(const interval& x, const T& y) return interval(rnd.div_down(xl, y), rnd.div_up(xu, y), true); } +template inline +BOOST_GPU_ENABLED interval operator/(const T2& x, const interval& y) +{ return static_cast(x) / y; } + } // namespace numeric } // namespace boost diff --git a/include/boost/numeric/interval/arith2.hpp b/include/boost/numeric/interval/arith2.hpp index eec87e0..459dc0c 100644 --- a/include/boost/numeric/interval/arith2.hpp +++ b/include/boost/numeric/interval/arith2.hpp @@ -29,8 +29,8 @@ namespace boost { namespace numeric { template inline -interval fmod(const interval& x, - const interval& y) +BOOST_GPU_ENABLED interval fmod(const interval& x, + const interval& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -42,7 +42,7 @@ interval fmod(const interval& x, } template inline -interval fmod(const interval& x, const T& y) +BOOST_GPU_ENABLED interval fmod(const interval& x, const T& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -53,7 +53,7 @@ interval fmod(const interval& x, const T& y) } template inline -interval fmod(const T& x, const interval& y) +BOOST_GPU_ENABLED interval fmod(const T& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) return interval::empty(); @@ -67,8 +67,8 @@ interval fmod(const T& x, const interval& y) namespace interval_lib { template inline -interval division_part1(const interval& x, - const interval& y, bool& b) +BOOST_GPU_ENABLED interval division_part1(const interval& x, + const interval& y, bool& b) { typedef interval I; b = false; @@ -90,15 +90,15 @@ interval division_part1(const interval& x, } template inline -interval division_part2(const interval& x, - const interval& y, bool b = true) +BOOST_GPU_ENABLED interval division_part2(const interval& x, + const interval& y, bool b = true) { if (!b) return interval::empty(); return detail::div_zero_part2(x, y); } template inline -interval multiplicative_inverse(const interval& x) +BOOST_GPU_ENABLED interval multiplicative_inverse(const interval& x) { typedef interval I; if (detail::test_input(x)) @@ -124,7 +124,7 @@ interval multiplicative_inverse(const interval& x) namespace detail { template inline -T pow_dn(const T& x_, int pwr, Rounding& rnd) // x and pwr are positive +BOOST_GPU_ENABLED T pow_dn(const T& x_, int pwr, Rounding& rnd) // x and pwr are positive { T x = x_; T y = (pwr & 1) ? x_ : static_cast(1); @@ -138,7 +138,7 @@ T pow_dn(const T& x_, int pwr, Rounding& rnd) // x and pwr are positive } template inline -T pow_up(const T& x_, int pwr, Rounding& rnd) // x and pwr are positive +BOOST_GPU_ENABLED T pow_up(const T& x_, int pwr, Rounding& rnd) // x and pwr are positive { T x = x_; T y = (pwr & 1) ? x_ : static_cast(1); @@ -155,7 +155,7 @@ T pow_up(const T& x_, int pwr, Rounding& rnd) // x and pwr are positive } // namespace interval_lib template inline -interval pow(const interval& x, int pwr) +BOOST_GPU_ENABLED interval pow(const interval& x, int pwr) { BOOST_USING_STD_MAX(); using interval_lib::detail::pow_dn; @@ -195,7 +195,7 @@ interval pow(const interval& x, int pwr) } template inline -interval sqrt(const interval& x) +BOOST_GPU_ENABLED interval sqrt(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x) || interval_lib::user::is_neg(x.upper())) @@ -206,7 +206,7 @@ interval sqrt(const interval& x) } template inline -interval square(const interval& x) +BOOST_GPU_ENABLED interval square(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -226,7 +226,7 @@ namespace interval_lib { namespace detail { template< class I > inline -I root_aux(typename I::base_type const &x, int k) // x and k are bigger than one +BOOST_GPU_ENABLED I root_aux(typename I::base_type const &x, int k) // x and k are bigger than one { typedef typename I::base_type T; T tk(k); @@ -240,7 +240,7 @@ I root_aux(typename I::base_type const &x, int k) // x and k are bigger than one } template< class I > inline // x is positive and k bigger than one -typename I::base_type root_aux_dn(typename I::base_type const &x, int k) +BOOST_GPU_ENABLED typename I::base_type root_aux_dn(typename I::base_type const &x, int k) { typedef typename I::base_type T; typedef typename I::traits_type Policies; @@ -252,7 +252,7 @@ typename I::base_type root_aux_dn(typename I::base_type const &x, int k) } template< class I > inline // x is positive and k bigger than one -typename I::base_type root_aux_up(typename I::base_type const &x, int k) +BOOST_GPU_ENABLED typename I::base_type root_aux_up(typename I::base_type const &x, int k) { typedef typename I::base_type T; typedef typename I::traits_type Policies; @@ -267,7 +267,7 @@ typename I::base_type root_aux_up(typename I::base_type const &x, int k) } // namespace interval_lib template< class T, class Policies > inline -interval nth_root(interval const &x, int k) +BOOST_GPU_ENABLED interval nth_root(interval const &x, int k) { typedef interval I; if (interval_lib::detail::test_input(x)) return I::empty(); diff --git a/include/boost/numeric/interval/arith3.hpp b/include/boost/numeric/interval/arith3.hpp index 518e618..da6efdd 100644 --- a/include/boost/numeric/interval/arith3.hpp +++ b/include/boost/numeric/interval/arith3.hpp @@ -23,7 +23,7 @@ namespace numeric { namespace interval_lib { template inline -I add(const typename I::base_type& x, const typename I::base_type& y) +BOOST_GPU_ENABLED I add(const typename I::base_type& x, const typename I::base_type& y) { typedef typename I::traits_type Policies; if (detail::test_input(x, y)) @@ -33,7 +33,7 @@ I add(const typename I::base_type& x, const typename I::base_type& y) } template inline -I sub(const typename I::base_type& x, const typename I::base_type& y) +BOOST_GPU_ENABLED I sub(const typename I::base_type& x, const typename I::base_type& y) { typedef typename I::traits_type Policies; if (detail::test_input(x, y)) @@ -43,7 +43,7 @@ I sub(const typename I::base_type& x, const typename I::base_type& y) } template inline -I mul(const typename I::base_type& x, const typename I::base_type& y) +BOOST_GPU_ENABLED I mul(const typename I::base_type& x, const typename I::base_type& y) { typedef typename I::traits_type Policies; if (detail::test_input(x, y)) @@ -53,7 +53,7 @@ I mul(const typename I::base_type& x, const typename I::base_type& y) } template inline -I div(const typename I::base_type& x, const typename I::base_type& y) +BOOST_GPU_ENABLED I div(const typename I::base_type& x, const typename I::base_type& y) { typedef typename I::traits_type Policies; if (detail::test_input(x, y) || user::is_zero(y)) diff --git a/include/boost/numeric/interval/checking.hpp b/include/boost/numeric/interval/checking.hpp index f40ba65..65e49be 100644 --- a/include/boost/numeric/interval/checking.hpp +++ b/include/boost/numeric/interval/checking.hpp @@ -14,6 +14,7 @@ #include #include #include +#include namespace boost { namespace numeric { @@ -21,53 +22,59 @@ namespace interval_lib { struct exception_create_empty { - void operator()() + BOOST_GPU_ENABLED void operator()() { - throw std::runtime_error("boost::interval: empty interval created"); + BOOST_NUMERIC_INTERVAL_throw("boost::interval: empty interval created"); } }; struct exception_invalid_number { - void operator()() + BOOST_GPU_ENABLED void operator()() { - throw std::invalid_argument("boost::interval: invalid number"); + BOOST_NUMERIC_INTERVAL_throw("boost::interval: invalid number"); } }; template struct checking_base { - static T pos_inf() + BOOST_GPU_ENABLED static T pos_inf() { - assert(std::numeric_limits::has_infinity); - return std::numeric_limits::infinity(); + BOOST_NUMERIC_INTERVAL_using_std(numeric_limits); + assert(numeric_limits::has_infinity); + return numeric_limits::infinity(); } - static T neg_inf() + BOOST_GPU_ENABLED static T neg_inf() { - assert(std::numeric_limits::has_infinity); - return -std::numeric_limits::infinity(); + BOOST_NUMERIC_INTERVAL_using_std(numeric_limits); + assert(numeric_limits::has_infinity); + return -numeric_limits::infinity(); } - static T nan() + BOOST_GPU_ENABLED static T nan() { - assert(std::numeric_limits::has_quiet_NaN); - return std::numeric_limits::quiet_NaN(); + BOOST_NUMERIC_INTERVAL_using_std(numeric_limits); + assert(numeric_limits::has_quiet_NaN); + return numeric_limits::quiet_NaN(); } - static bool is_nan(const T& x) + BOOST_GPU_ENABLED static bool is_nan(const T& x) { - return std::numeric_limits::has_quiet_NaN && (x != x); + BOOST_NUMERIC_INTERVAL_using_std(numeric_limits); + return numeric_limits::has_quiet_NaN && (x != x); } - static T empty_lower() + BOOST_GPU_ENABLED static T empty_lower() { - return (std::numeric_limits::has_quiet_NaN ? - std::numeric_limits::quiet_NaN() : static_cast(1)); + BOOST_NUMERIC_INTERVAL_using_std(numeric_limits); + return (numeric_limits::has_quiet_NaN) ? + numeric_limits::quiet_NaN() : static_cast(1); } - static T empty_upper() + BOOST_GPU_ENABLED static T empty_upper() { - return (std::numeric_limits::has_quiet_NaN ? - std::numeric_limits::quiet_NaN() : static_cast(0)); + BOOST_NUMERIC_INTERVAL_using_std(numeric_limits); + return (numeric_limits::has_quiet_NaN) ? + numeric_limits::quiet_NaN() : static_cast(0); } - static bool is_empty(const T& l, const T& u) + BOOST_GPU_ENABLED static bool is_empty(const T& l, const T& u) { return !(l <= u); // safety for partial orders } @@ -77,22 +84,22 @@ template, class Exception = exception_create_empty> struct checking_no_empty: Checking { - static T nan() + BOOST_GPU_ENABLED static T nan() { - assert(false); + assert(false); // ? return Checking::nan(); } - static T empty_lower() + BOOST_GPU_ENABLED static T empty_lower() { Exception()(); return Checking::empty_lower(); } - static T empty_upper() + BOOST_GPU_ENABLED static T empty_upper() { Exception()(); return Checking::empty_upper(); } - static bool is_empty(const T&, const T&) + BOOST_GPU_ENABLED static bool is_empty(const T&, const T&) { return false; } @@ -101,7 +108,7 @@ struct checking_no_empty: Checking template > struct checking_no_nan: Checking { - static bool is_nan(const T&) + BOOST_GPU_ENABLED static bool is_nan(const T&) { return false; } @@ -111,9 +118,11 @@ template, class Exception = exception_invalid_number> struct checking_catch_nan: Checking { - static bool is_nan(const T& x) + BOOST_GPU_ENABLED static bool is_nan(const T& x) { - if (Checking::is_nan(x)) Exception()(); + if (Checking::is_nan(x)) { + Exception()(); + } return false; } }; diff --git a/include/boost/numeric/interval/compare/certain.hpp b/include/boost/numeric/interval/compare/certain.hpp index 9232d5c..0b88217 100644 --- a/include/boost/numeric/interval/compare/certain.hpp +++ b/include/boost/numeric/interval/compare/certain.hpp @@ -20,86 +20,110 @@ namespace compare { namespace certain { template inline -bool operator<(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() < y.lower(); } template inline -bool operator<(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() < y; } template inline -bool operator<=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() <= y.lower(); } template inline -bool operator<=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() <= y; } template inline -bool operator>(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.lower() > y.upper(); } template inline -bool operator>(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.lower() > y; } template inline -bool operator>=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.lower() >= y.upper(); } template inline -bool operator>=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.lower() >= y; } template inline -bool operator==(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() == y.lower() && x.lower() == y.upper(); } template inline -bool operator==(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() == y && x.lower() == y; } template inline -bool operator!=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() < y.lower() || x.lower() > y.upper(); } template inline -bool operator!=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { + comparison_error()(); + } return x.upper() < y || x.lower() > y; } diff --git a/include/boost/numeric/interval/compare/explicit.hpp b/include/boost/numeric/interval/compare/explicit.hpp index f66dc64..6b6103d 100644 --- a/include/boost/numeric/interval/compare/explicit.hpp +++ b/include/boost/numeric/interval/compare/explicit.hpp @@ -22,109 +22,109 @@ namespace interval_lib { */ template inline -bool cerlt(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool cerlt(const interval& x, const interval& y) { return x.upper() < y.lower(); } template inline -bool cerlt(const interval& x, const T& y) +BOOST_GPU_ENABLED bool cerlt(const interval& x, const T& y) { return x.upper() < y; } template inline -bool cerlt(const T& x, const interval& y) +BOOST_GPU_ENABLED bool cerlt(const T& x, const interval& y) { return x < y.lower(); } template inline -bool cerle(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool cerle(const interval& x, const interval& y) { return x.upper() <= y.lower(); } template inline -bool cerle(const interval& x, const T& y) +BOOST_GPU_ENABLED bool cerle(const interval& x, const T& y) { return x.upper() <= y; } template inline -bool cerle(const T& x, const interval& y) +BOOST_GPU_ENABLED bool cerle(const T& x, const interval& y) { return x <= y.lower(); } template inline -bool cergt(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool cergt(const interval& x, const interval& y) { return x.lower() > y.upper(); } template inline -bool cergt(const interval& x, const T& y) +BOOST_GPU_ENABLED bool cergt(const interval& x, const T& y) { return x.lower() > y; } template inline -bool cergt(const T& x, const interval& y) +BOOST_GPU_ENABLED bool cergt(const T& x, const interval& y) { return x > y.upper(); } template inline -bool cerge(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool cerge(const interval& x, const interval& y) { return x.lower() >= y.upper(); } template inline -bool cerge(const interval& x, const T& y) +BOOST_GPU_ENABLED bool cerge(const interval& x, const T& y) { return x.lower() >= y; } template inline -bool cerge(const T& x, const interval& y) +BOOST_GPU_ENABLED bool cerge(const T& x, const interval& y) { return x >= y.upper(); } template inline -bool cereq(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool cereq(const interval& x, const interval& y) { return x.lower() == y.upper() && y.lower() == x.upper(); } template inline -bool cereq(const interval& x, const T& y) +BOOST_GPU_ENABLED bool cereq(const interval& x, const T& y) { return x.lower() == y && x.upper() == y; } template inline -bool cereq(const T& x, const interval& y) +BOOST_GPU_ENABLED bool cereq(const T& x, const interval& y) { return x == y.lower() && x == y.upper(); } template inline -bool cerne(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool cerne(const interval& x, const interval& y) { return x.upper() < y.lower() || y.upper() < x.lower(); } template inline -bool cerne(const interval& x, const T& y) +BOOST_GPU_ENABLED bool cerne(const interval& x, const T& y) { return x.upper() < y || y < x.lower(); } template inline -bool cerne(const T& x, const interval& y) +BOOST_GPU_ENABLED bool cerne(const T& x, const interval& y) { return x < y.lower() || y.upper() < x; } @@ -134,109 +134,109 @@ bool cerne(const T& x, const interval& y) */ template inline -bool poslt(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool poslt(const interval& x, const interval& y) { return x.lower() < y.upper(); } template inline -bool poslt(const interval& x, const T& y) +BOOST_GPU_ENABLED bool poslt(const interval& x, const T& y) { return x.lower() < y; } template inline -bool poslt(const T& x, const interval& y) +BOOST_GPU_ENABLED bool poslt(const T& x, const interval& y) { return x < y.upper(); } template inline -bool posle(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool posle(const interval& x, const interval& y) { return x.lower() <= y.upper(); } template inline -bool posle(const interval& x, const T& y) +BOOST_GPU_ENABLED bool posle(const interval& x, const T& y) { return x.lower() <= y; } template inline -bool posle(const T& x, const interval& y) +BOOST_GPU_ENABLED bool posle(const T& x, const interval& y) { return x <= y.upper(); } template inline -bool posgt(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool posgt(const interval& x, const interval& y) { return x.upper() > y.lower(); } template inline -bool posgt(const interval& x, const T& y) +BOOST_GPU_ENABLED bool posgt(const interval& x, const T& y) { return x.upper() > y; } template inline -bool posgt(const T& x, const interval & y) +BOOST_GPU_ENABLED bool posgt(const T& x, const interval & y) { return x > y.lower(); } template inline -bool posge(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool posge(const interval& x, const interval& y) { return x.upper() >= y.lower(); } template inline -bool posge(const interval& x, const T& y) +BOOST_GPU_ENABLED bool posge(const interval& x, const T& y) { return x.upper() >= y; } template inline -bool posge(const T& x, const interval& y) +BOOST_GPU_ENABLED bool posge(const T& x, const interval& y) { return x >= y.lower(); } template inline -bool poseq(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool poseq(const interval& x, const interval& y) { return x.upper() >= y.lower() && y.upper() >= x.lower(); } template inline -bool poseq(const interval& x, const T& y) +BOOST_GPU_ENABLED bool poseq(const interval& x, const T& y) { return x.upper() >= y && y >= x.lower(); } template inline -bool poseq(const T& x, const interval& y) +BOOST_GPU_ENABLED bool poseq(const T& x, const interval& y) { return x >= y.lower() && y.upper() >= x; } template inline -bool posne(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool posne(const interval& x, const interval& y) { return x.upper() != y.lower() || y.upper() != x.lower(); } template inline -bool posne(const interval& x, const T& y) +BOOST_GPU_ENABLED bool posne(const interval& x, const T& y) { return x.upper() != y || y != x.lower(); } template inline -bool posne(const T& x, const interval& y) +BOOST_GPU_ENABLED bool posne(const T& x, const interval& y) { return x != y.lower() || y.upper() != x; } diff --git a/include/boost/numeric/interval/compare/lexicographic.hpp b/include/boost/numeric/interval/compare/lexicographic.hpp index 03f6036..d421f69 100644 --- a/include/boost/numeric/interval/compare/lexicographic.hpp +++ b/include/boost/numeric/interval/compare/lexicographic.hpp @@ -20,96 +20,96 @@ namespace compare { namespace lexicographic { template inline -bool operator<(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } const T& xl = x.lower(); const T& yl = y.lower(); return xl < yl || (xl == yl && x.upper() < y.upper()); } template inline -bool operator<(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() < y; } template inline -bool operator<=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } const T& xl = x.lower(); const T& yl = y.lower(); return xl < yl || (xl == yl && x.upper() <= y.upper()); } template inline -bool operator<=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } const T& xl = x.lower(); return xl < y || (xl == y && x.upper() <= y); } template inline -bool operator>(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } const T& xl = x.lower(); const T& yl = y.lower(); return xl > yl || (xl == yl && x.upper() > y.upper()); } template inline -bool operator>(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } const T& xl = x.lower(); return xl > y || (xl == y && x.upper() > y); } template inline -bool operator>=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } const T& xl = x.lower(); const T& yl = y.lower(); return xl > yl || (xl == yl && x.upper() >= y.upper()); } template inline -bool operator>=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() >= y; } template inline -bool operator==(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() == y.lower() && x.upper() == y.upper(); } template inline -bool operator==(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() == y && x.upper() == y; } template inline -bool operator!=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() != y.lower() || x.upper() != y.upper(); } template inline -bool operator!=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() != y || x.upper() != y; } diff --git a/include/boost/numeric/interval/compare/possible.hpp b/include/boost/numeric/interval/compare/possible.hpp index 59bec31..98c1d5a 100644 --- a/include/boost/numeric/interval/compare/possible.hpp +++ b/include/boost/numeric/interval/compare/possible.hpp @@ -20,86 +20,86 @@ namespace compare { namespace possible { template inline -bool operator<(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() < y.upper(); } template inline -bool operator<(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() < y; } template inline -bool operator<=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() <= y.upper(); } template inline -bool operator<=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() <= y; } template inline -bool operator>(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.upper() > y.lower(); } template inline -bool operator>(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.upper() > y; } template inline -bool operator>=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.upper() >= y.lower(); } template inline -bool operator>=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.upper() >= y; } template inline -bool operator==(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() <= y.upper() && x.upper() >= y.lower(); } template inline -bool operator==(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() <= y && x.upper() >= y; } template inline -bool operator!=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const interval& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() != y.upper() || x.upper() != y.lower(); } template inline -bool operator!=(const interval& x, const T& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const T& y) { - if (detail::test_input(x, y)) throw comparison_error(); + if (detail::test_input(x, y)) { comparison_error()(); } return x.lower() != y || x.upper() != y; } diff --git a/include/boost/numeric/interval/compare/set.hpp b/include/boost/numeric/interval/compare/set.hpp index 44f4d6e..9674611 100644 --- a/include/boost/numeric/interval/compare/set.hpp +++ b/include/boost/numeric/interval/compare/set.hpp @@ -21,75 +21,81 @@ namespace compare { namespace set { template inline -bool operator<(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<(const interval& x, const interval& y) { return proper_subset(x, y); } template inline -bool operator<(const interval& , const T& ) +BOOST_GPU_ENABLED bool operator<(const interval& , const T& ) { - throw comparison_error(); + comparison_error()(); + return false; } template inline -bool operator<=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator<=(const interval& x, const interval& y) { return subset(x, y); } template inline -bool operator<=(const interval& , const T& ) +BOOST_GPU_ENABLED bool operator<=(const interval& , const T& ) { - throw comparison_error(); + comparison_error()(); + return false; } template inline -bool operator>(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>(const interval& x, const interval& y) { return proper_subset(y, x); } template inline -bool operator>(const interval& , const T& ) +BOOST_GPU_ENABLED bool operator>(const interval& , const T& ) { - throw comparison_error(); + comparison_error()(); + return false; } template inline -bool operator>=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator>=(const interval& x, const interval& y) { return subset(y, x); } template inline -bool operator>=(const interval& , const T& ) +BOOST_GPU_ENABLED bool operator>=(const interval& , const T& ) { - throw comparison_error(); + comparison_error()(); + return false; } template inline -bool operator==(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator==(const interval& x, const interval& y) { return equal(y, x); } template inline -bool operator==(const interval& , const T& ) +BOOST_GPU_ENABLED bool operator==(const interval& , const T& ) { - throw comparison_error(); + comparison_error()(); + return false; } template inline -bool operator!=(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool operator!=(const interval& x, const interval& y) { return !equal(y, x); } template inline -bool operator!=(const interval& , const T& ) +BOOST_GPU_ENABLED bool operator!=(const interval& , const T& ) { - throw comparison_error(); + comparison_error()(); + return false; } } // namespace set diff --git a/include/boost/numeric/interval/constants.hpp b/include/boost/numeric/interval/constants.hpp index 6a22834..6a2318b 100644 --- a/include/boost/numeric/interval/constants.hpp +++ b/include/boost/numeric/interval/constants.hpp @@ -24,38 +24,38 @@ static const float pi_f_u = 13176795.0f/(1<<22); static const double pi_d_l = (3373259426.0 + 273688.0 / (1<<21)) / (1<<30); static const double pi_d_u = (3373259426.0 + 273689.0 / (1<<21)) / (1<<30); -template inline T pi_lower() { return 3; } -template inline T pi_upper() { return 4; } -template inline T pi_half_lower() { return 1; } -template inline T pi_half_upper() { return 2; } -template inline T pi_twice_lower() { return 6; } -template inline T pi_twice_upper() { return 7; } +template inline BOOST_GPU_ENABLED T pi_lower() { return 3; } +template inline BOOST_GPU_ENABLED T pi_upper() { return 4; } +template inline BOOST_GPU_ENABLED T pi_half_lower() { return 1; } +template inline BOOST_GPU_ENABLED T pi_half_upper() { return 2; } +template inline BOOST_GPU_ENABLED T pi_twice_lower() { return 6; } +template inline BOOST_GPU_ENABLED T pi_twice_upper() { return 7; } -template<> inline float pi_lower() { return pi_f_l; } -template<> inline float pi_upper() { return pi_f_u; } -template<> inline float pi_half_lower() { return pi_f_l / 2; } -template<> inline float pi_half_upper() { return pi_f_u / 2; } -template<> inline float pi_twice_lower() { return pi_f_l * 2; } -template<> inline float pi_twice_upper() { return pi_f_u * 2; } +template<> inline BOOST_GPU_ENABLED float pi_lower() { return pi_f_l; } +template<> inline BOOST_GPU_ENABLED float pi_upper() { return pi_f_u; } +template<> inline BOOST_GPU_ENABLED float pi_half_lower() { return pi_f_l / 2; } +template<> inline BOOST_GPU_ENABLED float pi_half_upper() { return pi_f_u / 2; } +template<> inline BOOST_GPU_ENABLED float pi_twice_lower() { return pi_f_l * 2; } +template<> inline BOOST_GPU_ENABLED float pi_twice_upper() { return pi_f_u * 2; } -template<> inline double pi_lower() { return pi_d_l; } -template<> inline double pi_upper() { return pi_d_u; } -template<> inline double pi_half_lower() { return pi_d_l / 2; } -template<> inline double pi_half_upper() { return pi_d_u / 2; } -template<> inline double pi_twice_lower() { return pi_d_l * 2; } -template<> inline double pi_twice_upper() { return pi_d_u * 2; } +template<> inline BOOST_GPU_ENABLED double pi_lower() { return pi_d_l; } +template<> inline BOOST_GPU_ENABLED double pi_upper() { return pi_d_u; } +template<> inline BOOST_GPU_ENABLED double pi_half_lower() { return pi_d_l / 2; } +template<> inline BOOST_GPU_ENABLED double pi_half_upper() { return pi_d_u / 2; } +template<> inline BOOST_GPU_ENABLED double pi_twice_lower() { return pi_d_l * 2; } +template<> inline BOOST_GPU_ENABLED double pi_twice_upper() { return pi_d_u * 2; } -template<> inline long double pi_lower() { return pi_d_l; } -template<> inline long double pi_upper() { return pi_d_u; } -template<> inline long double pi_half_lower() { return pi_d_l / 2; } -template<> inline long double pi_half_upper() { return pi_d_u / 2; } -template<> inline long double pi_twice_lower() { return pi_d_l * 2; } -template<> inline long double pi_twice_upper() { return pi_d_u * 2; } +template<> inline BOOST_GPU_ENABLED long double pi_lower() { return pi_d_l; } +template<> inline BOOST_GPU_ENABLED long double pi_upper() { return pi_d_u; } +template<> inline BOOST_GPU_ENABLED long double pi_half_lower() { return pi_d_l / 2; } +template<> inline BOOST_GPU_ENABLED long double pi_half_upper() { return pi_d_u / 2; } +template<> inline BOOST_GPU_ENABLED long double pi_twice_lower() { return pi_d_l * 2; } +template<> inline BOOST_GPU_ENABLED long double pi_twice_upper() { return pi_d_u * 2; } } // namespace constants template inline -I pi() +BOOST_GPU_ENABLED I pi() { typedef typename I::base_type T; return I(constants::pi_lower(), @@ -63,7 +63,7 @@ I pi() } template inline -I pi_half() +BOOST_GPU_ENABLED I pi_half() { typedef typename I::base_type T; return I(constants::pi_half_lower(), @@ -71,7 +71,7 @@ I pi_half() } template inline -I pi_twice() +BOOST_GPU_ENABLED I pi_twice() { typedef typename I::base_type T; return I(constants::pi_twice_lower(), diff --git a/include/boost/numeric/interval/detail/bugs.hpp b/include/boost/numeric/interval/detail/bugs.hpp index e15bd0c..5107a20 100644 --- a/include/boost/numeric/interval/detail/bugs.hpp +++ b/include/boost/numeric/interval/detail/bugs.hpp @@ -45,4 +45,24 @@ # define BOOST_NUMERIC_INTERVAL_using_ahyp(a) #endif +#if defined(__CUDACC__) +# define BOOST_GPU_DISABLED __host__ +#else +# define BOOST_GPU_DISABLED +#endif + +#if defined(__CUDA_ARCH__) +# define BOOST_NUMERIC_INTERVAL_std(a) cuda::std::a +# undef BOOST_USING_STD_MIN +# define BOOST_USING_STD_MIN() using boost::numeric::gpu_spec::min +# undef BOOST_USING_STD_MAX +# define BOOST_USING_STD_MAX() using boost::numeric::gpu_spec::max +# define BOOST_NUMERIC_INTERVAL_throw(exception) printf("%s\n", exception); __trap() +#else +# define BOOST_NUMERIC_INTERVAL_std(a) std::a +# define BOOST_NUMERIC_INTERVAL_throw(exception) throw std::runtime_error(exception) +#endif + +#define BOOST_NUMERIC_INTERVAL_using_std(a) using BOOST_NUMERIC_INTERVAL_std(a) + #endif // BOOST_NUMERIC_INTERVAL_DETAIL_BUGS diff --git a/include/boost/numeric/interval/detail/c99sub_rounding_control.hpp b/include/boost/numeric/interval/detail/c99sub_rounding_control.hpp index a8684df..5a5f66f 100644 --- a/include/boost/numeric/interval/detail/c99sub_rounding_control.hpp +++ b/include/boost/numeric/interval/detail/c99sub_rounding_control.hpp @@ -40,4 +40,4 @@ struct c99_rounding_control } // namespace numeric } // namespace boost -#endif // BOOST_NUMERIC_INTERVAL_DETAIL_C99SUB_ROUBDING_CONTROL_HPP +#endif // BOOST_NUMERIC_INTERVAL_DETAIL_C99SUB_ROUNDING_CONTROL_HPP diff --git a/include/boost/numeric/interval/detail/cuda_rounding_control.hpp b/include/boost/numeric/interval/detail/cuda_rounding_control.hpp new file mode 100644 index 0000000..1a9f901 --- /dev/null +++ b/include/boost/numeric/interval/detail/cuda_rounding_control.hpp @@ -0,0 +1,265 @@ +/* Boost interval/detail/cuda_rounding_control.hpp file + * + * Copyright 2024 Lorenz Gillner + * + * Distributed under 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) + */ + +#ifndef BOOST_NUMERIC_INTERVAL_DETAIL_CUDA_ROUNDING_CONTROL_HPP +#define BOOST_NUMERIC_INTERVAL_DETAIL_CUDA_ROUNDING_CONTROL_HPP + +#if !defined(__CUDACC__) && !defined(__NVCC__) +#error Boost.Numeric.Interval: This header is intended for CUDA GPUs only. +#endif + +#include +#include + +#include +#include +#include +#include + +#include +#include + +namespace boost { +namespace numeric { +namespace gpu_spec { + + /** templates */ + template __device__ T min(const T &x, const T &y); + template __device__ T max(const T &x, const T &y); + + /** float specialization */ + template <> inline __device__ float min(const float &x, const float &y) + { return fminf(x, y); } + template <> inline __device__ float max(const float &x, const float &y) + { return fmaxf(x, y); } + + /** double specialization */ + template <> inline __device__ double min(const double &x, const double &y) + { return fmin(x, y); } + template <> inline __device__ double max(const double &x, const double &y) + { return fmax(x, y); } + +} // namespace gpu_spec + +namespace interval_lib { + +namespace detail { + + /** directed operation templates */ + template __device__ T add_rd(const T &x, const T &y); + template __device__ T add_ru(const T &x, const T &y); + template __device__ T sub_rd(const T &x, const T &y); + template __device__ T sub_ru(const T &x, const T &y); + template __device__ T mul_rd(const T &x, const T &y); + template __device__ T mul_ru(const T &x, const T &y); + template __device__ T div_rd(const T &x, const T &y); + template __device__ T div_ru(const T &x, const T &y); + template __device__ T sqrt_rd(const T &x); + template __device__ T sqrt_ru(const T &x); + template __device__ T pred(const T &x); + template __device__ T succ(const T &x); + template __device__ T mid(const T &x, const T &y); + + template __device__ T conv_rd(U const &v); + template __device__ T conv_ru(U const &v); + + /** float specialization */ + #define BOOST_NUMERIC_INTERVAL_gpu_spec(a) \ + template <> inline __device__ float a##_rd(const float &x, const float &y) \ + { return __f##a##_rd(x, y); } \ + template <> inline __device__ float a##_ru(const float &x, const float &y) \ + { return __f##a##_ru(x, y); } + BOOST_NUMERIC_INTERVAL_gpu_spec(add); + BOOST_NUMERIC_INTERVAL_gpu_spec(sub); + BOOST_NUMERIC_INTERVAL_gpu_spec(mul); + BOOST_NUMERIC_INTERVAL_gpu_spec(div); + #undef BOOST_NUMERIC_INTERVAL_gpu_spec + template <> inline __device__ float sqrt_rd(const float &x) + { return __fsqrt_rd(x); } + template <> inline __device__ float sqrt_ru(const float &x) + { return __fsqrt_ru(x); } + template <> inline __device__ float pred(const float &x) + { return nextafterf(x, -CUDART_INF_F); } + template <> inline __device__ float succ(const float &x) + { return nextafterf(x, CUDART_INF_F); } + template <> inline __device__ float mid(const float &x, const float &y) + { return __fdiv_rn(__fadd_rn(x, y), 2); } + template <> inline __device__ float conv_rd(const int &v) + { return __int2float_rd(v); } + template <> inline __device__ float conv_ru(const int &v) + { return __int2float_ru(v); } + template <> inline __device__ float conv_rd(const long long int &v) + { return __ll2float_rd(v); } + template <> inline __device__ float conv_ru(const long long int &v) + { return __ll2float_ru(v); } + template <> inline __device__ float conv_rd(const double &v) + { return __double2float_rd(v); } + template <> inline __device__ float conv_ru(const double &v) + { return __double2float_ru(v); } + + /** double specialization */ + #define BOOST_NUMERIC_INTERVAL_gpu_spec(a) \ + template <> inline __device__ double a##_rd(const double &x, const double &y) \ + { return __d##a##_rd(x, y); } \ + template <> inline __device__ double a##_ru(const double &x, const double &y) \ + { return __d##a##_ru(x, y); } + BOOST_NUMERIC_INTERVAL_gpu_spec(add); + BOOST_NUMERIC_INTERVAL_gpu_spec(sub); + BOOST_NUMERIC_INTERVAL_gpu_spec(mul); + BOOST_NUMERIC_INTERVAL_gpu_spec(div); + #undef BOOST_NUMERIC_INTERVAL_gpu_spec + template <> inline __device__ double sqrt_rd(const double &x) + { return __dsqrt_rd(x); } + template <> inline __device__ double sqrt_ru(const double &x) + { return __dsqrt_ru(x); } + template <> inline __device__ double pred(const double &x) + { return nextafter(x, -CUDART_INF); } + template <> inline __device__ double succ(const double &x) + { return nextafter(x, CUDART_INF); } + template <> inline __device__ double mid(const double &x, const double &y) + { return __ddiv_rn(__dadd_rn(x, y), 2); } + template <> inline __device__ double conv_rd(const int &v) + { return __int2double_rn(v); } + template <> inline __device__ double conv_ru(const int &v) + { return __int2double_rn(v); } + template <> inline __device__ double conv_rd(const long long int &v) + { return __ll2double_rd(v); } + template <> inline __device__ double conv_ru(const long long int &v) + { return __ll2double_ru(v); } + +} // namespace detail + +/** + * Checking should be compatible with the GPU "as is", so no further spezializations are needed + */ + +template +struct checking_base_gpu: checking_base +{}; + +/** + * GPU rounding control + */ + +template +struct rounding_control_gpu: rounding_control +{}; + +template <> +struct rounding_control_gpu +{ + __device__ float to_int(const float& x) { return nearbyintf(x); } +}; + +template <> +struct rounding_control_gpu +{ + __device__ double to_int(const double& x) { return nearbyint(x); } +}; + +/** + * Directly rounded arithmetic + */ + +template > +struct rounded_arith_direct; + +template +struct rounded_arith_direct : Rounding +{ + __device__ void init() {} + template __device__ T conv_down(U const &v) { return detail::conv_rd(v); } + template __device__ T conv_up(U const &v) { return detail::conv_ru(v); } + #define BOOST_NUMERIC_INTERVAL_new_func(a) \ + __device__ T a##_down(const T &x, const T &y) \ + { return detail::a##_rd(x, y); } \ + __device__ T a##_up(const T &x, const T &y) \ + { return detail::a##_ru(x, y); } + BOOST_NUMERIC_INTERVAL_new_func(add); + BOOST_NUMERIC_INTERVAL_new_func(sub); + BOOST_NUMERIC_INTERVAL_new_func(mul); + BOOST_NUMERIC_INTERVAL_new_func(div); + #undef BOOST_NUMERIC_INTERVAL_new_func + __device__ T sqrt_down(const T &x) { return detail::sqrt_rd(x); } + __device__ T sqrt_up(const T &x) { return detail::sqrt_ru(x); } + __device__ T median(const T &x, const T &y) { return detail::mid(x, y); } + __device__ T int_down(const T &x) { BOOST_NUMERIC_INTERVAL_using_math(floor); return floor(x); } + __device__ T int_up(const T &x) { BOOST_NUMERIC_INTERVAL_using_math(ceil); return ceil(x); } + __device__ T next_down(const T &x) { return detail::pred(x); } + __device__ T next_up(const T &x) { return detail::succ(x); } +}; + +/** + * Rounded transcendental functions + */ + +template > +struct rounded_transc_direct; + +template +struct rounded_transc_direct: Rounding +{ + #define BOOST_NUMERIC_INTERVAL_new_func(a) \ + __device__ T a##_down(const T& x) \ + { BOOST_NUMERIC_INTERVAL_using_math(a); T ax = a(x); return (ax > T(-1) && ax != T(0)) ? next_down(ax) : ax; } \ + __device__ T a##_up (const T& x) \ + { BOOST_NUMERIC_INTERVAL_using_math(a); T ax = a(x); return (ax < T(1) && ax != T(0)) ? next_up(ax) : ax; } + BOOST_NUMERIC_INTERVAL_new_func(exp) // 1 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(log) // 1 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(sin) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(cos) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(tan) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(asin) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(acos) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(atan) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(sinh) // 2 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(cosh) // 1 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(tanh) // 1 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(asinh) // 3 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(acosh) // 3 ulp (double) + BOOST_NUMERIC_INTERVAL_new_func(atanh) // 2 ulp (double) + #undef BOOST_NUMERIC_INTERVAL_new_func +}; + +/** + * Convenience wrappers for GPU use + */ + +template +struct rounded_arith_gpu: rounded_arith_direct +{}; + +template +struct rounded_transc_gpu: rounded_transc_direct +{}; + +/* fall back to exact mode if no specialization is available */ +template +struct rounded_math_gpu: save_state_nothing > +{}; + +template <> +struct rounded_math_gpu: save_state_nothing > +{}; + +template <> +struct rounded_math_gpu: save_state_nothing > +{}; + +template +struct default_policies_gpu +{ + typedef policies, checking_base_gpu > type; +}; + +} // namespace interval_lib +} // namespace numeric +} // namespace boost + +#endif /* BOOST_NUMERIC_INTERVAL_DETAIL_CUDA_ROUNDING_CONTROL_HPP */ \ No newline at end of file diff --git a/include/boost/numeric/interval/detail/division.hpp b/include/boost/numeric/interval/detail/division.hpp index 24fb025..0382278 100644 --- a/include/boost/numeric/interval/detail/division.hpp +++ b/include/boost/numeric/interval/detail/division.hpp @@ -22,8 +22,8 @@ namespace interval_lib { namespace detail { template inline -interval div_non_zero(const interval& x, - const interval& y) +BOOST_GPU_ENABLED interval div_non_zero(const interval& x, + const interval& y) { // assert(!in_zero(y)); typename Policies::rounding rnd; @@ -50,7 +50,7 @@ interval div_non_zero(const interval& x, } template inline -interval div_non_zero(const T& x, const interval& y) +BOOST_GPU_ENABLED interval div_non_zero(const T& x, const interval& y) { // assert(!in_zero(y)); typename Policies::rounding rnd; @@ -64,7 +64,7 @@ interval div_non_zero(const T& x, const interval& y) } template inline -interval div_positive(const interval& x, const T& yu) +BOOST_GPU_ENABLED interval div_positive(const interval& x, const T& yu) { // assert(::boost::numeric::interval_lib::user::is_pos(yu)); if (::boost::numeric::interval_lib::user::is_zero(x.lower()) && @@ -84,7 +84,7 @@ interval div_positive(const interval& x, const T& yu) } template inline -interval div_positive(const T& x, const T& yu) +BOOST_GPU_ENABLED interval div_positive(const T& x, const T& yu) { // assert(::boost::numeric::interval_lib::user::is_pos(yu)); typedef interval I; @@ -99,7 +99,7 @@ interval div_positive(const T& x, const T& yu) } template inline -interval div_negative(const interval& x, const T& yl) +BOOST_GPU_ENABLED interval div_negative(const interval& x, const T& yl) { // assert(::boost::numeric::interval_lib::user::is_neg(yl)); if (::boost::numeric::interval_lib::user::is_zero(x.lower()) && @@ -119,7 +119,7 @@ interval div_negative(const interval& x, const T& yl) } template inline -interval div_negative(const T& x, const T& yl) +BOOST_GPU_ENABLED interval div_negative(const T& x, const T& yl) { // assert(::boost::numeric::interval_lib::user::is_neg(yl)); typedef interval I; @@ -134,7 +134,7 @@ interval div_negative(const T& x, const T& yl) } template inline -interval div_zero(const interval& x) +BOOST_GPU_ENABLED interval div_zero(const interval& x) { if (::boost::numeric::interval_lib::user::is_zero(x.lower()) && ::boost::numeric::interval_lib::user::is_zero(x.upper())) @@ -143,7 +143,7 @@ interval div_zero(const interval& x) } template inline -interval div_zero(const T& x) +BOOST_GPU_ENABLED interval div_zero(const T& x) { if (::boost::numeric::interval_lib::user::is_zero(x)) return interval(static_cast(0), static_cast(0), true); @@ -151,7 +151,7 @@ interval div_zero(const T& x) } template inline -interval div_zero_part1(const interval& x, +BOOST_GPU_ENABLED interval div_zero_part1(const interval& x, const interval& y, bool& b) { // assert(::boost::numeric::interval_lib::user::is_neg(y.lower()) && ::boost::numeric::interval_lib::user::is_pos(y.upper())); @@ -173,8 +173,8 @@ interval div_zero_part1(const interval& x, } template inline -interval div_zero_part2(const interval& x, - const interval& y) +BOOST_GPU_ENABLED interval div_zero_part2(const interval& x, + const interval& y) { // assert(::boost::numeric::interval_lib::user::is_neg(y.lower()) && ::boost::numeric::interval_lib::user::is_pos(y.upper()) && (div_zero_part1(x, y, b), b)); typename Policies::rounding rnd; diff --git a/include/boost/numeric/interval/detail/interval_prototype.hpp b/include/boost/numeric/interval/detail/interval_prototype.hpp index 47ba447..8f9734c 100644 --- a/include/boost/numeric/interval/detail/interval_prototype.hpp +++ b/include/boost/numeric/interval/detail/interval_prototype.hpp @@ -10,6 +10,8 @@ #ifndef BOOST_NUMERIC_INTERVAL_DETAIL_INTERVAL_PROTOTYPE_HPP #define BOOST_NUMERIC_INTERVAL_DETAIL_INTERVAL_PROTOTYPE_HPP +#include + namespace boost { namespace numeric { diff --git a/include/boost/numeric/interval/detail/test_input.hpp b/include/boost/numeric/interval/detail/test_input.hpp index 9f93bb0..11b70d3 100644 --- a/include/boost/numeric/interval/detail/test_input.hpp +++ b/include/boost/numeric/interval/detail/test_input.hpp @@ -17,53 +17,53 @@ namespace numeric { namespace interval_lib { namespace user { -template inline -bool is_zero(T const &v) { return v == static_cast(0); } +template +BOOST_GPU_ENABLED bool is_zero(T const &v) { return v == static_cast(0); } -template inline -bool is_neg (T const &v) { return v < static_cast(0); } +template +BOOST_GPU_ENABLED bool is_neg (T const &v) { return v < static_cast(0); } -template inline -bool is_pos (T const &v) { return v > static_cast(0); } +template +BOOST_GPU_ENABLED bool is_pos (T const &v) { return v > static_cast(0); } } // namespace user namespace detail { -template inline -bool test_input(const interval& x) { +template +BOOST_GPU_ENABLED bool test_input(const interval& x) { typedef typename Policies::checking checking; return checking::is_empty(x.lower(), x.upper()); } -template inline -bool test_input(const interval& x, const interval& y) { +template +BOOST_GPU_ENABLED bool test_input(const interval& x, const interval& y) { typedef typename Policies1::checking checking1; typedef typename Policies2::checking checking2; return checking1::is_empty(x.lower(), x.upper()) || checking2::is_empty(y.lower(), y.upper()); } -template inline -bool test_input(const T& x, const interval& y) { +template +BOOST_GPU_ENABLED bool test_input(const T& x, const interval& y) { typedef typename Policies::checking checking; return checking::is_nan(x) || checking::is_empty(y.lower(), y.upper()); } -template inline -bool test_input(const interval& x, const T& y) { +template +BOOST_GPU_ENABLED bool test_input(const interval& x, const T& y) { typedef typename Policies::checking checking; return checking::is_empty(x.lower(), x.upper()) || checking::is_nan(y); } -template inline -bool test_input(const T& x) { +template +BOOST_GPU_ENABLED bool test_input(const T& x) { typedef typename Policies::checking checking; return checking::is_nan(x); } -template inline -bool test_input(const T& x, const T& y) { +template +BOOST_GPU_ENABLED bool test_input(const T& x, const T& y) { typedef typename Policies::checking checking; return checking::is_nan(x) || checking::is_nan(y); } diff --git a/include/boost/numeric/interval/ext/integer.hpp b/include/boost/numeric/interval/ext/integer.hpp index 628a343..8669180 100644 --- a/include/boost/numeric/interval/ext/integer.hpp +++ b/include/boost/numeric/interval/ext/integer.hpp @@ -17,49 +17,49 @@ namespace boost { namespace numeric { template inline -interval operator+ (const interval& x, int y) +BOOST_GPU_ENABLED interval operator+ (const interval& x, int y) { return x + static_cast(y); } template inline -interval operator+ (int x, const interval& y) +BOOST_GPU_ENABLED interval operator+ (int x, const interval& y) { return static_cast(x) + y; } template inline -interval operator- (const interval& x, int y) +BOOST_GPU_ENABLED interval operator- (const interval& x, int y) { return x - static_cast(y); } template inline -interval operator- (int x, const interval& y) +BOOST_GPU_ENABLED interval operator- (int x, const interval& y) { return static_cast(x) - y; } template inline -interval operator* (const interval& x, int y) +BOOST_GPU_ENABLED interval operator* (const interval& x, int y) { return x * static_cast(y); } template inline -interval operator* (int x, const interval& y) +BOOST_GPU_ENABLED interval operator* (int x, const interval& y) { return static_cast(x) * y; } template inline -interval operator/ (const interval& x, int y) +BOOST_GPU_ENABLED interval operator/ (const interval& x, int y) { return x / static_cast(y); } template inline -interval operator/ (int x, const interval& y) +BOOST_GPU_ENABLED interval operator/ (int x, const interval& y) { return static_cast(x) / y; } diff --git a/include/boost/numeric/interval/hw_rounding.hpp b/include/boost/numeric/interval/hw_rounding.hpp index 46d452e..4b10887 100644 --- a/include/boost/numeric/interval/hw_rounding.hpp +++ b/include/boost/numeric/interval/hw_rounding.hpp @@ -34,6 +34,10 @@ # include #endif +#if defined(__CUDACC__) || defined(__NVCC__) +# include +#endif + #if defined(BOOST_NUMERIC_INTERVAL_NO_HARDWARE) && !defined(BOOST_NO_FENV_H) # include #endif diff --git a/include/boost/numeric/interval/interval.hpp b/include/boost/numeric/interval/interval.hpp index e89af48..ef54c00 100644 --- a/include/boost/numeric/interval/interval.hpp +++ b/include/boost/numeric/interval/interval.hpp @@ -18,14 +18,13 @@ namespace boost { namespace numeric { namespace interval_lib { - -class comparison_error - : public std::runtime_error + +struct comparison_error { -public: - comparison_error() - : std::runtime_error("boost::interval: uncertain comparison") - { } + BOOST_GPU_ENABLED void operator()() + { + BOOST_NUMERIC_INTERVAL_throw("boost::interval: uncertain comparison"); + } }; } // namespace interval_lib @@ -44,69 +43,76 @@ class interval typedef T base_type; typedef Policies traits_type; - T const &lower() const; - T const &upper() const; - - interval(); - interval(T const &v); - template interval(T1 const &v); - interval(T const &l, T const &u); - template interval(T1 const &l, T2 const &u); - interval(interval const &r); - template interval(interval const &r); - template interval(interval const &r); - - interval &operator=(T const &v); - template interval &operator=(T1 const &v); - interval &operator=(interval const &r); - template interval &operator=(interval const &r); - template interval &operator=(interval const &r); - - void assign(const T& l, const T& u); - - static interval empty(); - static interval whole(); - static interval hull(const T& x, const T& y); - - interval& operator+= (const T& r); - interval& operator+= (const interval& r); - interval& operator-= (const T& r); - interval& operator-= (const interval& r); - interval& operator*= (const T& r); - interval& operator*= (const interval& r); - interval& operator/= (const T& r); - interval& operator/= (const interval& r); - - bool operator< (const interval_holder& r) const; - bool operator> (const interval_holder& r) const; - bool operator<= (const interval_holder& r) const; - bool operator>= (const interval_holder& r) const; - bool operator== (const interval_holder& r) const; - bool operator!= (const interval_holder& r) const; - - bool operator< (const number_holder& r) const; - bool operator> (const number_holder& r) const; - bool operator<= (const number_holder& r) const; - bool operator>= (const number_holder& r) const; - bool operator== (const number_holder& r) const; - bool operator!= (const number_holder& r) const; + BOOST_GPU_ENABLED T const &lower() const; + BOOST_GPU_ENABLED T const &upper() const; + + BOOST_GPU_ENABLED interval(); + BOOST_GPU_ENABLED interval(T const &v); + template + BOOST_GPU_ENABLED interval(T1 const &v); + BOOST_GPU_ENABLED interval(T const &l, T const &u); + template + BOOST_GPU_ENABLED interval(T1 const &l, T2 const &u); + BOOST_GPU_ENABLED interval(interval const &r); + template + BOOST_GPU_ENABLED interval(interval const &r); + template + BOOST_GPU_ENABLED interval(interval const &r); + + BOOST_GPU_ENABLED interval &operator=(T const &v); + template + BOOST_GPU_ENABLED interval &operator=(T1 const &v); + BOOST_GPU_ENABLED interval &operator=(interval const &r); + template + BOOST_GPU_ENABLED interval &operator=(interval const &r); + template + BOOST_GPU_ENABLED interval &operator=(interval const &r); + + BOOST_GPU_ENABLED void assign(const T& l, const T& u); + + BOOST_GPU_ENABLED static interval empty(); + BOOST_GPU_ENABLED static interval whole(); + BOOST_GPU_ENABLED static interval hull(const T& x, const T& y); + + BOOST_GPU_ENABLED interval& operator+= (const T& r); + BOOST_GPU_ENABLED interval& operator+= (const interval& r); + BOOST_GPU_ENABLED interval& operator-= (const T& r); + BOOST_GPU_ENABLED interval& operator-= (const interval& r); + BOOST_GPU_ENABLED interval& operator*= (const T& r); + BOOST_GPU_ENABLED interval& operator*= (const interval& r); + BOOST_GPU_ENABLED interval& operator/= (const T& r); + BOOST_GPU_ENABLED interval& operator/= (const interval& r); + + BOOST_GPU_ENABLED bool operator< (const interval_holder& r) const; + BOOST_GPU_ENABLED bool operator> (const interval_holder& r) const; + BOOST_GPU_ENABLED bool operator<= (const interval_holder& r) const; + BOOST_GPU_ENABLED bool operator>= (const interval_holder& r) const; + BOOST_GPU_ENABLED bool operator== (const interval_holder& r) const; + BOOST_GPU_ENABLED bool operator!= (const interval_holder& r) const; + + BOOST_GPU_ENABLED bool operator< (const number_holder& r) const; + BOOST_GPU_ENABLED bool operator> (const number_holder& r) const; + BOOST_GPU_ENABLED bool operator<= (const number_holder& r) const; + BOOST_GPU_ENABLED bool operator>= (const number_holder& r) const; + BOOST_GPU_ENABLED bool operator== (const number_holder& r) const; + BOOST_GPU_ENABLED bool operator!= (const number_holder& r) const; // the following is for internal use only, it is not a published interface // nevertheless, it's public because friends don't always work correctly. - interval(const T& l, const T& u, bool): low(l), up(u) {} - void set_empty(); - void set_whole(); - void set(const T& l, const T& u); + BOOST_GPU_ENABLED interval(const T& l, const T& u, bool): low(l), up(u) {} + BOOST_GPU_ENABLED void set_empty(); + BOOST_GPU_ENABLED void set_whole(); + BOOST_GPU_ENABLED void set(const T& l, const T& u); private: struct interval_holder { template - interval_holder(const interval& r) + BOOST_GPU_ENABLED interval_holder(const interval& r) : low(r.lower()), up(r.upper()) { typedef typename Policies2::checking checking2; if (checking2::is_empty(low, up)) - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); } const T& low; @@ -114,11 +120,11 @@ class interval }; struct number_holder { - number_holder(const T& r) : val(r) + BOOST_GPU_ENABLED number_holder(const T& r) : val(r) { typedef typename Policies::checking checking; if (checking::is_nan(r)) - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); } const T& val; @@ -132,18 +138,18 @@ class interval }; template inline -interval::interval(): +BOOST_GPU_ENABLED interval::interval(): low(static_cast(0)), up(static_cast(0)) {} template inline -interval::interval(T const &v): low(v), up(v) +BOOST_GPU_ENABLED interval::interval(T const &v): low(v), up(v) { if (checking::is_nan(v)) set_empty(); } template template inline -interval::interval(T1 const &v) +BOOST_GPU_ENABLED interval::interval(T1 const &v) { if (checking::is_nan(v)) set_empty(); else { @@ -154,7 +160,7 @@ interval::interval(T1 const &v) } template template inline -interval::interval(T1 const &l, T2 const &u) +BOOST_GPU_ENABLED interval::interval(T1 const &l, T2 const &u) { if (checking::is_nan(l) || checking::is_nan(u) || !(l <= u)) set_empty(); else { @@ -165,7 +171,7 @@ interval::interval(T1 const &l, T2 const &u) } template inline -interval::interval(T const &l, T const &u): low(l), up(u) +BOOST_GPU_ENABLED interval::interval(T const &l, T const &u): low(l), up(u) { if (checking::is_nan(l) || checking::is_nan(u) || !(l <= u)) set_empty(); @@ -173,18 +179,18 @@ interval::interval(T const &l, T const &u): low(l), up(u) template inline -interval::interval(interval const &r): low(r.lower()), up(r.upper()) +BOOST_GPU_ENABLED interval::interval(interval const &r): low(r.lower()), up(r.upper()) {} template template inline -interval::interval(interval const &r): low(r.lower()), up(r.upper()) +BOOST_GPU_ENABLED interval::interval(interval const &r): low(r.lower()), up(r.upper()) { typedef typename Policies1::checking checking1; if (checking1::is_empty(r.lower(), r.upper())) set_empty(); } template template inline -interval::interval(interval const &r) +BOOST_GPU_ENABLED interval::interval(interval const &r) { typedef typename Policies1::checking checking1; if (checking1::is_empty(r.lower(), r.upper())) set_empty(); @@ -196,7 +202,7 @@ interval::interval(interval const &r) } template inline -interval &interval::operator=(T const &v) +BOOST_GPU_ENABLED interval &interval::operator=(T const &v) { if (checking::is_nan(v)) set_empty(); else low = up = v; @@ -204,7 +210,7 @@ interval &interval::operator=(T const &v) } template template inline -interval &interval::operator=(T1 const &v) +BOOST_GPU_ENABLED interval &interval::operator=(T1 const &v) { if (checking::is_nan(v)) set_empty(); else { @@ -216,7 +222,7 @@ interval &interval::operator=(T1 const &v) } template inline -interval &interval::operator=(interval const &r) +BOOST_GPU_ENABLED interval &interval::operator=(interval const &r) { low = r.lower(); up = r.upper(); @@ -224,7 +230,7 @@ interval &interval::operator=(interval co } template template inline -interval &interval::operator=(interval const &r) +BOOST_GPU_ENABLED interval &interval::operator=(interval const &r) { typedef typename Policies1::checking checking1; if (checking1::is_empty(r.lower(), r.upper())) set_empty(); @@ -236,7 +242,7 @@ interval &interval::operator=(interval c } template template inline -interval &interval::operator=(interval const &r) +BOOST_GPU_ENABLED interval &interval::operator=(interval const &r) { typedef typename Policies1::checking checking1; if (checking1::is_empty(r.lower(), r.upper())) set_empty(); @@ -249,7 +255,7 @@ interval &interval::operator=(interval } template inline -void interval::assign(const T& l, const T& u) +BOOST_GPU_ENABLED void interval::assign(const T& l, const T& u) { if (checking::is_nan(l) || checking::is_nan(u) || !(l <= u)) set_empty(); @@ -257,28 +263,28 @@ void interval::assign(const T& l, const T& u) } template inline -void interval::set(const T& l, const T& u) +BOOST_GPU_ENABLED void interval::set(const T& l, const T& u) { low = l; up = u; } template inline -void interval::set_empty() +BOOST_GPU_ENABLED void interval::set_empty() { low = checking::empty_lower(); up = checking::empty_upper(); } template inline -void interval::set_whole() +BOOST_GPU_ENABLED void interval::set_whole() { low = checking::neg_inf(); up = checking::pos_inf(); } template inline -interval interval::hull(const T& x, const T& y) +BOOST_GPU_ENABLED interval interval::hull(const T& x, const T& y) { bool bad_x = checking::is_nan(x); bool bad_y = checking::is_nan(y); @@ -292,26 +298,26 @@ interval interval::hull(const T& x, const T& y) } template inline -interval interval::empty() +BOOST_GPU_ENABLED interval interval::empty() { return interval(checking::empty_lower(), checking::empty_upper(), true); } template inline -interval interval::whole() +BOOST_GPU_ENABLED interval interval::whole() { return interval(checking::neg_inf(), checking::pos_inf(), true); } template inline -const T& interval::lower() const +BOOST_GPU_ENABLED const T& interval::lower() const { return low; } template inline -const T& interval::upper() const +BOOST_GPU_ENABLED const T& interval::upper() const { return up; } @@ -321,63 +327,69 @@ const T& interval::upper() const */ template inline -bool interval::operator< (const interval_holder& r) const +BOOST_GPU_ENABLED bool interval::operator< (const interval_holder& r) const { if (!checking::is_empty(low, up)) { if (up < r.low) return true; else if (low >= r.up) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator> (const interval_holder& r) const +BOOST_GPU_ENABLED bool interval::operator> (const interval_holder& r) const { if (!checking::is_empty(low, up)) { if (low > r.up) return true; else if (up <= r.low) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator<= (const interval_holder& r) const +BOOST_GPU_ENABLED bool interval::operator<= (const interval_holder& r) const { if (!checking::is_empty(low, up)) { if (up <= r.low) return true; else if (low > r.up) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator>= (const interval_holder& r) const +BOOST_GPU_ENABLED bool interval::operator>= (const interval_holder& r) const { if (!checking::is_empty(low, up)) { if (low >= r.up) return true; else if (up < r.low) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator== (const interval_holder& r) const +BOOST_GPU_ENABLED bool interval::operator== (const interval_holder& r) const { if (!checking::is_empty(low, up)) { if (up == r.low && low == r.up) return true; else if (up < r.low || low > r.up) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator!= (const interval_holder& r) const +BOOST_GPU_ENABLED bool interval::operator!= (const interval_holder& r) const { if (!checking::is_empty(low, up)) { if (up < r.low || low > r.up) return true; else if (up == r.low && low == r.up) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } /* @@ -385,63 +397,69 @@ bool interval::operator!= (const interval_holder& r) const */ template inline -bool interval::operator< (const number_holder& r) const +BOOST_GPU_ENABLED bool interval::operator< (const number_holder& r) const { if (!checking::is_empty(low, up)) { if (up < r.val) return true; else if (low >= r.val) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator> (const number_holder& r) const +BOOST_GPU_ENABLED bool interval::operator> (const number_holder& r) const { if (!checking::is_empty(low, up)) { if (low > r.val) return true; else if (up <= r.val) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator<= (const number_holder& r) const +BOOST_GPU_ENABLED bool interval::operator<= (const number_holder& r) const { if (!checking::is_empty(low, up)) { if (up <= r.val) return true; else if (low > r.val) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator>= (const number_holder& r) const +BOOST_GPU_ENABLED bool interval::operator>= (const number_holder& r) const { if (!checking::is_empty(low, up)) { if (low >= r.val) return true; else if (up < r.val) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator== (const number_holder& r) const +BOOST_GPU_ENABLED bool interval::operator== (const number_holder& r) const { if (!checking::is_empty(low, up)) { if (up == r.val && low == r.val) return true; else if (up < r.val || low > r.val) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } template inline -bool interval::operator!= (const number_holder& r) const +BOOST_GPU_ENABLED bool interval::operator!= (const number_holder& r) const { if (!checking::is_empty(low, up)) { if (up < r.val || low > r.val) return true; else if (up == r.val && low == r.val) return false; } - throw interval_lib::comparison_error(); + interval_lib::comparison_error()(); + return false; // never reached } } // namespace numeric diff --git a/include/boost/numeric/interval/limits.hpp b/include/boost/numeric/interval/limits.hpp index ce080dd..75b863a 100644 --- a/include/boost/numeric/interval/limits.hpp +++ b/include/boost/numeric/interval/limits.hpp @@ -26,18 +26,18 @@ class numeric_limits > typedef boost::numeric::interval I; typedef numeric_limits bl; public: - static I min BOOST_PREVENT_MACRO_SUBSTITUTION () BOOST_NOEXCEPT_OR_NOTHROW { return I((bl::min)(), (bl::min)()); } - static I max BOOST_PREVENT_MACRO_SUBSTITUTION () BOOST_NOEXCEPT_OR_NOTHROW { return I((bl::max)(), (bl::max)()); } - static I epsilon() BOOST_NOEXCEPT_OR_NOTHROW { return I(bl::epsilon(), bl::epsilon()); } + static BOOST_GPU_ENABLED I min BOOST_PREVENT_MACRO_SUBSTITUTION () BOOST_NOEXCEPT_OR_NOTHROW { return I((bl::min)(), (bl::min)()); } + static BOOST_GPU_ENABLED I max BOOST_PREVENT_MACRO_SUBSTITUTION () BOOST_NOEXCEPT_OR_NOTHROW { return I((bl::max)(), (bl::max)()); } + static BOOST_GPU_ENABLED I epsilon() BOOST_NOEXCEPT_OR_NOTHROW { return I(bl::epsilon(), bl::epsilon()); } BOOST_STATIC_CONSTANT(float_round_style, round_style = round_indeterminate); BOOST_STATIC_CONSTANT(bool, is_iec559 = false); - static I infinity () BOOST_NOEXCEPT_OR_NOTHROW { return I::whole(); } - static I quiet_NaN() BOOST_NOEXCEPT_OR_NOTHROW { return I::empty(); } - static I signaling_NaN() BOOST_NOEXCEPT_OR_NOTHROW + static BOOST_GPU_ENABLED I infinity () BOOST_NOEXCEPT_OR_NOTHROW { return I::whole(); } + static BOOST_GPU_ENABLED I quiet_NaN() BOOST_NOEXCEPT_OR_NOTHROW { return I::empty(); } + static BOOST_GPU_ENABLED I signaling_NaN() BOOST_NOEXCEPT_OR_NOTHROW { return I(bl::signaling_NaN(), bl::signaling_Nan()); } - static I denorm_min() BOOST_NOEXCEPT_OR_NOTHROW + static BOOST_GPU_ENABLED I denorm_min() BOOST_NOEXCEPT_OR_NOTHROW { return I(bl::denorm_min(), bl::denorm_min()); } private: static I round_error(); // hide this on purpose, not yet implemented diff --git a/include/boost/numeric/interval/rounded_arith.hpp b/include/boost/numeric/interval/rounded_arith.hpp index 4bb8691..78e7a6d 100644 --- a/include/boost/numeric/interval/rounded_arith.hpp +++ b/include/boost/numeric/interval/rounded_arith.hpp @@ -25,25 +25,27 @@ namespace interval_lib { template struct rounded_arith_exact: Rounding { - void init() { } - template T conv_down(U const &v) { return v; } - template T conv_up (U const &v) { return v; } - T add_down (const T& x, const T& y) { return x + y; } - T add_up (const T& x, const T& y) { return x + y; } - T sub_down (const T& x, const T& y) { return x - y; } - T sub_up (const T& x, const T& y) { return x - y; } - T mul_down (const T& x, const T& y) { return x * y; } - T mul_up (const T& x, const T& y) { return x * y; } - T div_down (const T& x, const T& y) { return x / y; } - T div_up (const T& x, const T& y) { return x / y; } - T median (const T& x, const T& y) { return (x + y) / 2; } - T sqrt_down(const T& x) + BOOST_GPU_ENABLED void init() { } + template + BOOST_GPU_ENABLED T conv_down(U const &v) { return v; } + template + BOOST_GPU_ENABLED T conv_up (U const &v) { return v; } + BOOST_GPU_ENABLED T add_down (const T& x, const T& y) { return x + y; } + BOOST_GPU_ENABLED T add_up (const T& x, const T& y) { return x + y; } + BOOST_GPU_ENABLED T sub_down (const T& x, const T& y) { return x - y; } + BOOST_GPU_ENABLED T sub_up (const T& x, const T& y) { return x - y; } + BOOST_GPU_ENABLED T mul_down (const T& x, const T& y) { return x * y; } + BOOST_GPU_ENABLED T mul_up (const T& x, const T& y) { return x * y; } + BOOST_GPU_ENABLED T div_down (const T& x, const T& y) { return x / y; } + BOOST_GPU_ENABLED T div_up (const T& x, const T& y) { return x / y; } + BOOST_GPU_ENABLED T median (const T& x, const T& y) { return (x + y) / 2; } + BOOST_GPU_ENABLED T sqrt_down(const T& x) { BOOST_NUMERIC_INTERVAL_using_math(sqrt); return sqrt(x); } - T sqrt_up (const T& x) + BOOST_GPU_ENABLED T sqrt_up (const T& x) { BOOST_NUMERIC_INTERVAL_using_math(sqrt); return sqrt(x); } - T int_down (const T& x) + BOOST_GPU_ENABLED T int_down (const T& x) { BOOST_NUMERIC_INTERVAL_using_math(floor); return floor(x); } - T int_up (const T& x) + BOOST_GPU_ENABLED T int_up (const T& x) { BOOST_NUMERIC_INTERVAL_using_math(ceil); return ceil(x); } }; diff --git a/include/boost/numeric/interval/rounded_transc.hpp b/include/boost/numeric/interval/rounded_transc.hpp index ef6e28b..8ac90b6 100644 --- a/include/boost/numeric/interval/rounded_transc.hpp +++ b/include/boost/numeric/interval/rounded_transc.hpp @@ -22,8 +22,8 @@ template struct rounded_transc_exact: Rounding { # define BOOST_NUMERIC_INTERVAL_new_func(f) \ - T f##_down(const T& x) { BOOST_NUMERIC_INTERVAL_using_math(f); return f(x); } \ - T f##_up (const T& x) { BOOST_NUMERIC_INTERVAL_using_math(f); return f(x); } + BOOST_GPU_ENABLED T f##_down(const T& x) { BOOST_NUMERIC_INTERVAL_using_math(f); return f(x); } \ + BOOST_GPU_ENABLED T f##_up (const T& x) { BOOST_NUMERIC_INTERVAL_using_math(f); return f(x); } BOOST_NUMERIC_INTERVAL_new_func(exp) BOOST_NUMERIC_INTERVAL_new_func(log) BOOST_NUMERIC_INTERVAL_new_func(sin) @@ -37,8 +37,8 @@ struct rounded_transc_exact: Rounding BOOST_NUMERIC_INTERVAL_new_func(tanh) # undef BOOST_NUMERIC_INTERVAL_new_func # define BOOST_NUMERIC_INTERVAL_new_func(f) \ - T f##_down(const T& x) { BOOST_NUMERIC_INTERVAL_using_ahyp(f); return f(x); } \ - T f##_up (const T& x) { BOOST_NUMERIC_INTERVAL_using_ahyp(f); return f(x); } + BOOST_GPU_ENABLED T f##_down(const T& x) { BOOST_NUMERIC_INTERVAL_using_ahyp(f); return f(x); } \ + BOOST_GPU_ENABLED T f##_up (const T& x) { BOOST_NUMERIC_INTERVAL_using_ahyp(f); return f(x); } BOOST_NUMERIC_INTERVAL_new_func(asinh) BOOST_NUMERIC_INTERVAL_new_func(acosh) BOOST_NUMERIC_INTERVAL_new_func(atanh) diff --git a/include/boost/numeric/interval/rounding.hpp b/include/boost/numeric/interval/rounding.hpp index bf5b3c2..4dfd0ff 100644 --- a/include/boost/numeric/interval/rounding.hpp +++ b/include/boost/numeric/interval/rounding.hpp @@ -22,13 +22,13 @@ template struct rounding_control { typedef int rounding_mode; - static void get_rounding_mode(rounding_mode&) {} - static void set_rounding_mode(rounding_mode) {} - static void upward() {} - static void downward() {} - static void to_nearest() {} - static const T& to_int(const T& x) { return x; } - static const T& force_rounding(const T& x) { return x; } + static BOOST_GPU_ENABLED void get_rounding_mode(rounding_mode&) {} + static BOOST_GPU_ENABLED void set_rounding_mode(rounding_mode) {} + static BOOST_GPU_ENABLED void upward() {} + static BOOST_GPU_ENABLED void downward() {} + static BOOST_GPU_ENABLED void to_nearest() {} + static BOOST_GPU_ENABLED const T& to_int(const T& x) { return x; } + static BOOST_GPU_ENABLED const T& force_rounding(const T& x) { return x; } }; /* diff --git a/include/boost/numeric/interval/transc.hpp b/include/boost/numeric/interval/transc.hpp index b9b7e60..0539fd0 100644 --- a/include/boost/numeric/interval/transc.hpp +++ b/include/boost/numeric/interval/transc.hpp @@ -25,7 +25,7 @@ namespace boost { namespace numeric { template inline -interval exp(const interval& x) +BOOST_GPU_ENABLED interval exp(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -35,7 +35,7 @@ interval exp(const interval& x) } template inline -interval log(const interval& x) +BOOST_GPU_ENABLED interval log(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x) || @@ -49,7 +49,7 @@ interval log(const interval& x) } template inline -interval cos(const interval& x) +BOOST_GPU_ENABLED interval cos(const interval& x) { if (interval_lib::detail::test_input(x)) return interval::empty(); @@ -78,7 +78,7 @@ interval cos(const interval& x) } template inline -interval sin(const interval& x) +BOOST_GPU_ENABLED interval sin(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -91,7 +91,7 @@ interval sin(const interval& x) } template inline -interval tan(const interval& x) +BOOST_GPU_ENABLED interval tan(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -111,7 +111,7 @@ interval tan(const interval& x) } template inline -interval asin(const interval& x) +BOOST_GPU_ENABLED interval asin(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x) @@ -128,7 +128,7 @@ interval asin(const interval& x) } template inline -interval acos(const interval& x) +BOOST_GPU_ENABLED interval acos(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x) @@ -145,7 +145,7 @@ interval acos(const interval& x) } template inline -interval atan(const interval& x) +BOOST_GPU_ENABLED interval atan(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -155,7 +155,7 @@ interval atan(const interval& x) } template inline -interval sinh(const interval& x) +BOOST_GPU_ENABLED interval sinh(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -165,7 +165,7 @@ interval sinh(const interval& x) } template inline -interval cosh(const interval& x) +BOOST_GPU_ENABLED interval cosh(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -180,7 +180,7 @@ interval cosh(const interval& x) } template inline -interval tanh(const interval& x) +BOOST_GPU_ENABLED interval tanh(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -190,7 +190,7 @@ interval tanh(const interval& x) } template inline -interval asinh(const interval& x) +BOOST_GPU_ENABLED interval asinh(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x)) @@ -200,7 +200,7 @@ interval asinh(const interval& x) } template inline -interval acosh(const interval& x) +BOOST_GPU_ENABLED interval acosh(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x) || x.upper() < static_cast(1)) @@ -211,7 +211,7 @@ interval acosh(const interval& x) } template inline -interval atanh(const interval& x) +BOOST_GPU_ENABLED interval atanh(const interval& x) { typedef interval I; if (interval_lib::detail::test_input(x) diff --git a/include/boost/numeric/interval/utility.hpp b/include/boost/numeric/interval/utility.hpp index 8e784b8..70f3060 100644 --- a/include/boost/numeric/interval/utility.hpp +++ b/include/boost/numeric/interval/utility.hpp @@ -13,7 +13,6 @@ #include #include -#include #include #include @@ -29,19 +28,19 @@ namespace numeric { */ template inline -const T& lower(const interval& x) +BOOST_GPU_ENABLED const T& lower(const interval& x) { return x.lower(); } template inline -const T& upper(const interval& x) +BOOST_GPU_ENABLED const T& upper(const interval& x) { return x.upper(); } template inline -T checked_lower(const interval& x) +BOOST_GPU_ENABLED T checked_lower(const interval& x) { if (empty(x)) { typedef typename Policies::checking checking; @@ -51,7 +50,7 @@ T checked_lower(const interval& x) } template inline -T checked_upper(const interval& x) +BOOST_GPU_ENABLED T checked_upper(const interval& x) { if (empty(x)) { typedef typename Policies::checking checking; @@ -61,7 +60,7 @@ T checked_upper(const interval& x) } template inline -T width(const interval& x) +BOOST_GPU_ENABLED T width(const interval& x) { if (interval_lib::detail::test_input(x)) return static_cast(0); typename Policies::rounding rnd; @@ -69,7 +68,7 @@ T width(const interval& x) } template inline -T median(const interval& x) +BOOST_GPU_ENABLED T median(const interval& x) { if (interval_lib::detail::test_input(x)) { typedef typename Policies::checking checking; @@ -80,7 +79,7 @@ T median(const interval& x) } template inline -interval widen(const interval& x, const T& v) +BOOST_GPU_ENABLED interval widen(const interval& x, const T& v) { if (interval_lib::detail::test_input(x)) return interval::empty(); @@ -94,34 +93,34 @@ interval widen(const interval& x, const T& v) */ template inline -bool empty(const interval& x) +BOOST_GPU_ENABLED bool empty(const interval& x) { return interval_lib::detail::test_input(x); } template inline -bool zero_in(const interval& x) +BOOST_GPU_ENABLED bool zero_in(const interval& x) { if (interval_lib::detail::test_input(x)) return false; return (!interval_lib::user::is_pos(x.lower())) && (!interval_lib::user::is_neg(x.upper())); } -template inline +/* template inline bool in_zero(const interval& x) // DEPRECATED { return zero_in(x); -} +} */ template inline -bool in(const T& x, const interval& y) +BOOST_GPU_ENABLED bool in(const T& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) return false; return y.lower() <= x && x <= y.upper(); } template inline -bool subset(const interval& x, +BOOST_GPU_ENABLED bool subset(const interval& x, const interval& y) { if (empty(x)) return true; @@ -129,7 +128,7 @@ bool subset(const interval& x, } template inline -bool proper_subset(const interval& x, +BOOST_GPU_ENABLED bool proper_subset(const interval& x, const interval& y) { if (empty(y)) return false; @@ -139,7 +138,7 @@ bool proper_subset(const interval& x, } template inline -bool overlap(const interval& x, +BOOST_GPU_ENABLED bool overlap(const interval& x, const interval& y) { if (interval_lib::detail::test_input(x, y)) return false; @@ -148,21 +147,21 @@ bool overlap(const interval& x, } template inline -bool singleton(const interval& x) +BOOST_GPU_ENABLED bool singleton(const interval& x) { return !empty(x) && x.lower() == x.upper(); } template inline -bool equal(const interval& x, const interval& y) +BOOST_GPU_ENABLED bool equal(const interval& x, const interval& y) { if (empty(x)) return empty(y); return !empty(y) && x.lower() == y.lower() && x.upper() == y.upper(); } template inline -interval intersect(const interval& x, - const interval& y) +BOOST_GPU_ENABLED interval intersect(const interval& x, + const interval& y) { BOOST_USING_STD_MIN(); BOOST_USING_STD_MAX(); @@ -175,7 +174,7 @@ interval intersect(const interval& x, } template inline -interval hull(const interval& x, +BOOST_GPU_ENABLED interval hull(const interval& x, const interval& y) { BOOST_USING_STD_MIN(); @@ -192,7 +191,7 @@ interval hull(const interval& x, } template inline -interval hull(const interval& x, const T& y) +BOOST_GPU_ENABLED interval hull(const interval& x, const T& y) { BOOST_USING_STD_MIN(); BOOST_USING_STD_MAX(); @@ -208,7 +207,7 @@ interval hull(const interval& x, const T& y) } template inline -interval hull(const T& x, const interval& y) +BOOST_GPU_ENABLED interval hull(const T& x, const interval& y) { BOOST_USING_STD_MIN(); BOOST_USING_STD_MAX(); @@ -224,20 +223,21 @@ interval hull(const T& x, const interval& y) } template inline -interval hull(const T& x, const T& y) +BOOST_GPU_ENABLED interval hull(const T& x, const T& y) { return interval::hull(x, y); } template inline -std::pair, interval > +BOOST_GPU_ENABLED BOOST_NUMERIC_INTERVAL_std(pair),interval > bisect(const interval& x) { + BOOST_NUMERIC_INTERVAL_using_std(pair); typedef interval I; if (interval_lib::detail::test_input(x)) - return std::pair(I::empty(), I::empty()); + return pair(I::empty(), I::empty()); const T m = median(x); - return std::pair(I(x.lower(), m, true), I(m, x.upper(), true)); + return pair(I(x.lower(), m, true), I(m, x.upper(), true)); } /* diff --git a/include/boost/numeric/interval/utility_fwd.hpp b/include/boost/numeric/interval/utility_fwd.hpp index 5fce06c..db57509 100644 --- a/include/boost/numeric/interval/utility_fwd.hpp +++ b/include/boost/numeric/interval/utility_fwd.hpp @@ -12,6 +12,7 @@ #include #include +#include #include /* @@ -25,94 +26,94 @@ namespace boost { namespace numeric { */ template - const T& lower(const interval& x); + BOOST_GPU_ENABLED const T& lower(const interval& x); template - const T& upper(const interval& x); + BOOST_GPU_ENABLED const T& upper(const interval& x); template - T checked_lower(const interval& x); + BOOST_GPU_ENABLED T checked_lower(const interval& x); template - T width(const interval& x); + BOOST_GPU_ENABLED T width(const interval& x); template - T median(const interval& x); + BOOST_GPU_ENABLED T median(const interval& x); template - interval widen(const interval& x, const T& v); + BOOST_GPU_ENABLED interval widen(const interval& x, const T& v); /* * Set-like operations */ template - bool empty(const interval& x); + BOOST_GPU_ENABLED bool empty(const interval& x); template - bool zero_in(const interval& x); + BOOST_GPU_ENABLED bool zero_in(const interval& x); template - bool in_zero(const interval& x); // DEPRECATED + BOOST_GPU_ENABLED bool in_zero(const interval& x); // DEPRECATED template - bool in(const T& x, const interval& y); + BOOST_GPU_ENABLED bool in(const T& x, const interval& y); template - bool + BOOST_GPU_ENABLED bool subset( const interval& x , const interval& y ); template - bool + BOOST_GPU_ENABLED bool proper_subset( const interval& x , const interval& y ); template - bool + BOOST_GPU_ENABLED bool overlap( const interval& x , const interval& y ); template - bool singleton(const interval& x); + BOOST_GPU_ENABLED bool singleton(const interval& x); template - bool + BOOST_GPU_ENABLED bool equal( const interval& x , const interval& y ); template - interval + BOOST_GPU_ENABLED interval intersect( const interval& x , const interval& y ); template - interval + BOOST_GPU_ENABLED interval hull(const interval& x, const interval& y); template - interval + BOOST_GPU_ENABLED interval hull(const interval& x, const T& y); template - interval + BOOST_GPU_ENABLED interval hull(const T& x, const interval& y); template - interval hull(const T& x, const T& y); - + BOOST_GPU_ENABLED interval hull(const T& x, const T& y); + template - std::pair,interval > + BOOST_GPU_ENABLED BOOST_NUMERIC_INTERVAL_std(pair),interval > bisect(const interval& x); /* @@ -120,48 +121,48 @@ namespace boost { namespace numeric { */ template - T norm(const interval& x); + BOOST_GPU_ENABLED T norm(const interval& x); template - interval abs(const interval& x); + BOOST_GPU_ENABLED interval abs(const interval& x); template - interval + BOOST_GPU_ENABLED interval max BOOST_PREVENT_MACRO_SUBSTITUTION ( const interval& x , const interval& y ); template - interval + BOOST_GPU_ENABLED interval max BOOST_PREVENT_MACRO_SUBSTITUTION ( const interval& x , const T& y ); template - interval + BOOST_GPU_ENABLED interval max BOOST_PREVENT_MACRO_SUBSTITUTION ( const T& x , const interval& y ); template - interval + BOOST_GPU_ENABLED interval min BOOST_PREVENT_MACRO_SUBSTITUTION ( const interval& x , const interval& y ); template - interval + BOOST_GPU_ENABLED interval min BOOST_PREVENT_MACRO_SUBSTITUTION ( const interval& x , const T& y ); template - interval + BOOST_GPU_ENABLED interval min BOOST_PREVENT_MACRO_SUBSTITUTION ( const T& x , const interval& y