From ef1ad8ac8cf69d5f068f08f46a498e4507e95e4f Mon Sep 17 00:00:00 2001 From: Allison Vacanti <alliepiper16@gmail.com> Date: Thu, 16 Jul 2020 18:36:53 -0400 Subject: [PATCH] Use transparent functionals in placeholder expressions. Fixes and adds regression tests for #1178 & #1229. --- testing/find.cu | 33 +++ testing/inner_product.cu | 18 ++ thrust/detail/functional/actor.h | 80 ++--- thrust/detail/functional/actor.inl | 158 +++------- .../operators/arithmetic_operators.h | 172 ++++++----- .../operators/assignment_operator.h | 22 +- .../functional/operators/bitwise_operators.h | 119 +++++--- .../operators/compound_assignment_operators.h | 275 ++++++++++++------ .../functional/operators/logical_operators.h | 28 +- .../functional/operators/operator_adaptors.h | 148 ++++++---- .../operators/relational_operators.h | 72 ++--- 11 files changed, 625 insertions(+), 500 deletions(-) diff --git a/testing/find.cu b/testing/find.cu index 427c8a7230..9252171dda 100644 --- a/testing/find.cu +++ b/testing/find.cu @@ -1,4 +1,5 @@ #include <unittest/unittest.h> +#include <thrust/sequence.h> #include <thrust/find.h> #include <thrust/iterator/retag.h> @@ -338,3 +339,35 @@ void TestFindWithBigIndexes() TestFindWithBigIndexesHelper(33); } DECLARE_UNITTEST(TestFindWithBigIndexes); + +namespace +{ + +class Weird +{ + int value; + +public: + __host__ __device__ Weird(int val, int) + : value(val) + {} + + friend __host__ __device__ + bool operator==(int x, Weird y) + { + return x == y.value; + } +}; + +} // end anon namespace + +void TestFindAsymmetricEquality() +{ // Regression test for thrust/thrust#1229 + thrust::host_vector<int> v(1000); + thrust::sequence(v.begin(), v.end()); + thrust::device_vector<int> dv(v); + auto result = thrust::find(dv.begin(), dv.end(), Weird(333, 0)); + ASSERT_EQUAL(*result, 333); + ASSERT_EQUAL(result - dv.begin(), 333); +} +DECLARE_UNITTEST(TestFindAsymmetricEquality); diff --git a/testing/inner_product.cu b/testing/inner_product.cu index 1bb897e6d5..07cce1dc17 100644 --- a/testing/inner_product.cu +++ b/testing/inner_product.cu @@ -1,8 +1,11 @@ #include <unittest/unittest.h> #include <thrust/inner_product.h> + +#include <thrust/functional.h> #include <thrust/iterator/retag.h> #include <thrust/device_malloc.h> #include <thrust/device_free.h> +#include <thrust/device_vector.h> template <class Vector> void TestInnerProductSimple(void) @@ -153,3 +156,18 @@ void TestInnerProductWithBigIndexes() TestInnerProductWithBigIndexesHelper(33); } DECLARE_UNITTEST(TestInnerProductWithBigIndexes); + +void TestInnerProductPlaceholders() +{ // Regression test for thrust/thrust#1178 + using namespace thrust::placeholders; + + thrust::device_vector<float> v1(100, 1.f); + thrust::device_vector<float> v2(100, 1.f); + + auto result = thrust::inner_product(v1.begin(), v1.end(), v2.begin(), 0.0f, + thrust::plus<float>{}, + _1 * _2 + 1.0f); + + ASSERT_ALMOST_EQUAL(result, 200.f); +} +DECLARE_UNITTEST(TestInnerProductPlaceholders); diff --git a/thrust/detail/functional/actor.h b/thrust/detail/functional/actor.h index 5759f79e38..63c32119f2 100644 --- a/thrust/detail/functional/actor.h +++ b/thrust/detail/functional/actor.h @@ -30,6 +30,7 @@ #include <thrust/detail/functional/value.h> #include <thrust/detail/functional/composite.h> #include <thrust/detail/functional/operators/assignment_operator.h> +#include <thrust/detail/raw_reference_cast.h> #include <thrust/detail/type_traits/result_of_adaptable_function.h> namespace thrust @@ -39,6 +40,34 @@ namespace detail namespace functional { +template <typename T, typename = void> +struct eval_ref_impl +{ + using type = T&; +}; + +template <typename T> +struct eval_ref_impl< + T, + typename std::enable_if<thrust::detail::is_wrapped_reference<T>::value>::type> +{ + using type = T; +}; + +// eval_ref<T> is +// - T when T is a subclass of thrust::reference +// - T& otherwise +// This is used to let thrust::references pass through actor evaluations. +template <typename T> +using eval_ref = + typename eval_ref_impl< + typename thrust::detail::remove_reference<T>::type + >::type; + +// Convenience def for tuple of eval_ref<T>s: +template <typename ...T> +using eval_tuple = thrust::tuple<eval_ref<T>...>; + template<typename Action, typename Env> struct apply_actor { @@ -61,55 +90,10 @@ template<typename Eval> typename apply_actor<eval_type, thrust::null_type >::type operator()(void) const; - template<typename T0> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&> >::type - operator()(T0 &_0) const; - - template<typename T0, typename T1> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&> >::type - operator()(T0 &_0, T1 &_1) const; - - template<typename T0, typename T1, typename T2> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2) const; - - template<typename T0, typename T1, typename T2, typename T3> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3) const; - - template<typename T0, typename T1, typename T2, typename T3, typename T4> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&,T4&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4) const; - - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5) const; - - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6) const; - - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7) const; - - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8> - __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8) const; - - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9> + template <typename... Ts> __host__ __device__ - typename apply_actor<eval_type, thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&,T9&> >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8, T9 &_9) const; + typename apply_actor<eval_type, eval_tuple<Ts...>>::type + operator()(Ts&&... ts) const; template<typename T> __host__ __device__ diff --git a/thrust/detail/functional/actor.inl b/thrust/detail/functional/actor.inl index 2c7fadd366..ac03e4a9cf 100644 --- a/thrust/detail/functional/actor.inl +++ b/thrust/detail/functional/actor.inl @@ -62,135 +62,43 @@ template<typename Eval> return eval_type::eval(thrust::null_type()); } // end basic_environment::operator() -template<typename Eval> - template<typename T0> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&> - >::type - actor<Eval> - ::operator()(T0 &_0) const -{ - return eval_type::eval(thrust::tie(_0)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1) const -{ - return eval_type::eval(thrust::tie(_0,_1)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3, typename T4> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&,T4&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5,_6)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5,_6,_7)); -} // end basic_environment::operator() - -template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8) const +struct actor_check_ref_type { - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5,_6,_7,_8)); -} // end basic_environment::operator() + // actor::operator() needs to construct a tuple of references to its + // arguments. To make this work with thrust::reference<T>, we need to + // detect thrust proxy references and store them as T rather than T&. + // This check ensures that the forwarding references passed into + // actor::operator() are either: + // - T&& if and only if T is a thrust::reference<U>, or + // - T& for any other types. + // This struct provides a nicer diagnostic for when these conditions aren't + // met. + template <typename T, typename... Ts> + __host__ __device__ + static void validate(T&& t, Ts&&... ts) + { + static constexpr bool is_lvalue_ref = std::is_lvalue_reference<T>::value; + static constexpr bool is_thrust_ref = + thrust::detail::is_wrapped_reference<T>::value; + static_assert(is_lvalue_ref || is_thrust_ref, + "Actor evaluations only support rvalue references to " + "thrust::reference subclasses."); + validate(ts...); + } + + __host__ __device__ static void validate() {} +}; template<typename Eval> - template<typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9> - __host__ __device__ - typename apply_actor< - typename actor<Eval>::eval_type, - typename thrust::tuple<T0&,T1&,T2&,T3&,T4&,T5&,T6&,T7&,T8&,T9&> - >::type - actor<Eval> - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8, T9 &_9) const +template<typename... Ts> +__host__ __device__ +typename apply_actor<typename actor<Eval>::eval_type, eval_tuple<Ts...>>::type +actor<Eval>::operator()(Ts&&... ts) const { - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5,_6,_7,_8,_9)); -} // end basic_environment::operator() + actor_check_ref_type::validate(ts...); + using tuple_type = eval_tuple<Ts...>; + return eval_type::eval(tuple_type(THRUST_FWD(ts)...)); +} // end actor<Eval>::operator() template<typename Eval> template<typename T> diff --git a/thrust/detail/functional/operators/arithmetic_operators.h b/thrust/detail/functional/operators/arithmetic_operators.h index 6628917d60..bd5b707e3b 100644 --- a/thrust/detail/functional/operators/arithmetic_operators.h +++ b/thrust/detail/functional/operators/arithmetic_operators.h @@ -33,49 +33,56 @@ template<typename Eval> __host__ __device__ actor< composite< - unary_operator<thrust::negate>, + transparent_unary_operator<thrust::negate<>>, actor<Eval> > > __host__ __device__ operator-(const actor<Eval> &_1) { - return compose(unary_operator<thrust::negate>(), _1); + return compose(transparent_unary_operator<thrust::negate<>>(), _1); } // end operator-() // there's no standard unary_plus functional, so roll an ad hoc one here -template<typename T> - struct unary_plus - : public thrust::unary_function<T,T> +struct unary_plus { - __host__ __device__ T operator()(const T &x) const {return +x;} -}; // end unary_plus + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1> + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(+THRUST_FWD(t1))) -> decltype(+THRUST_FWD(t1)) + { + return +THRUST_FWD(t1); + } +}; template<typename Eval> __host__ __device__ actor< composite< - unary_operator<unary_plus>, + transparent_unary_operator<unary_plus>, actor<Eval> > > operator+(const actor<Eval> &_1) { - return compose(unary_operator<unary_plus>(), _1); + return compose(transparent_unary_operator<unary_plus>(), _1); } // end operator+() template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::plus>, + transparent_binary_operator<thrust::plus<>>, actor<T1>, typename as_actor<T2>::type > > operator+(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::plus>(), + return compose(transparent_binary_operator<thrust::plus<>>(), make_actor(_1), make_actor(_2)); } // end operator+() @@ -84,14 +91,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::plus>, + transparent_binary_operator<thrust::plus<>>, typename as_actor<T1>::type, actor<T2> > > operator+(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::plus>(), + return compose(transparent_binary_operator<thrust::plus<>>(), make_actor(_1), make_actor(_2)); } // end operator+() @@ -100,14 +107,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::plus>, + transparent_binary_operator<thrust::plus<>>, actor<T1>, actor<T2> > > operator+(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::plus>(), + return compose(transparent_binary_operator<thrust::plus<>>(), make_actor(_1), make_actor(_2)); } // end operator+() @@ -116,14 +123,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::minus>, + transparent_binary_operator<thrust::minus<>>, typename as_actor<T1>::type, actor<T2> > > operator-(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::minus>(), + return compose(transparent_binary_operator<thrust::minus<>>(), make_actor(_1), make_actor(_2)); } // end operator-() @@ -132,14 +139,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::minus>, + transparent_binary_operator<thrust::minus<>>, actor<T1>, typename as_actor<T2>::type > > operator-(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::minus>(), + return compose(transparent_binary_operator<thrust::minus<>>(), make_actor(_1), make_actor(_2)); } // end operator-() @@ -148,14 +155,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::minus>, + transparent_binary_operator<thrust::minus<>>, actor<T1>, actor<T2> > > operator-(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::minus>(), + return compose(transparent_binary_operator<thrust::minus<>>(), make_actor(_1), make_actor(_2)); } // end operator-() @@ -164,14 +171,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::multiplies>, + transparent_binary_operator<thrust::multiplies<>>, typename as_actor<T1>::type, actor<T2> > > operator*(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::multiplies>(), + return compose(transparent_binary_operator<thrust::multiplies<>>(), make_actor(_1), make_actor(_2)); } // end operator*() @@ -180,14 +187,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::multiplies>, + transparent_binary_operator<thrust::multiplies<>>, actor<T1>, typename as_actor<T2>::type > > operator*(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::multiplies>(), + return compose(transparent_binary_operator<thrust::multiplies<>>(), make_actor(_1), make_actor(_2)); } // end operator*() @@ -196,14 +203,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::multiplies>, + transparent_binary_operator<thrust::multiplies<>>, actor<T1>, actor<T2> > > operator*(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::multiplies>(), + return compose(transparent_binary_operator<thrust::multiplies<>>(), make_actor(_1), make_actor(_2)); } // end operator*() @@ -212,14 +219,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::divides>, + transparent_binary_operator<thrust::divides<>>, actor<T1>, typename as_actor<T2>::type > > operator/(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::divides>(), + return compose(transparent_binary_operator<thrust::divides<>>(), make_actor(_1), make_actor(_2)); } // end operator/() @@ -228,14 +235,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::divides>, + transparent_binary_operator<thrust::divides<>>, typename as_actor<T1>::type, actor<T2> > > operator/(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::divides>(), + return compose(transparent_binary_operator<thrust::divides<>>(), make_actor(_1), make_actor(_2)); } // end operator/() @@ -244,14 +251,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::divides>, + transparent_binary_operator<thrust::divides<>>, actor<T1>, actor<T2> > > operator/(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::divides>(), + return compose(transparent_binary_operator<thrust::divides<>>(), make_actor(_1), make_actor(_2)); } // end operator/() @@ -260,14 +267,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::modulus>, + transparent_binary_operator<thrust::modulus<>>, actor<T1>, typename as_actor<T2>::type > > operator%(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::modulus>(), + return compose(transparent_binary_operator<thrust::modulus<>>(), make_actor(_1), make_actor(_2)); } // end operator%() @@ -276,14 +283,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::modulus>, + transparent_binary_operator<thrust::modulus<>>, typename as_actor<T1>::type, actor<T2> > > operator%(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::modulus>(), + return compose(transparent_binary_operator<thrust::modulus<void>>(), make_actor(_1), make_actor(_2)); } // end operator%() @@ -292,100 +299,131 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::modulus>, + transparent_binary_operator<thrust::modulus<>>, actor<T1>, actor<T2> > > operator%(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::modulus>(), + return compose(transparent_binary_operator<thrust::modulus<>>(), make_actor(_1), make_actor(_2)); } // end operator%() // there's no standard prefix_increment functional, so roll an ad hoc one here -template<typename T> - struct prefix_increment - : public thrust::unary_function<T&,T&> +struct prefix_increment { - __host__ __device__ T& operator()(T &x) const { return ++x; } + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1> + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(++THRUST_FWD(t1))) -> decltype(++THRUST_FWD(t1)) + { + return ++THRUST_FWD(t1); + } }; // end prefix_increment template<typename Eval> __host__ __device__ actor< composite< - unary_operator<prefix_increment>, + transparent_unary_operator<prefix_increment>, actor<Eval> > > operator++(const actor<Eval> &_1) { - return compose(unary_operator<prefix_increment>(), _1); + return compose(transparent_unary_operator<prefix_increment>(), _1); } // end operator++() -// there's no standard suffix_increment functional, so roll an ad hoc one here -template<typename T> - struct suffix_increment - : public thrust::unary_function<T&,T> + +// there's no standard postfix_increment functional, so roll an ad hoc one here +struct postfix_increment { - __host__ __device__ T operator()(T &x) const { return x++; } -}; // end suffix_increment + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1> + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(THRUST_FWD(t1)++)) -> decltype(THRUST_FWD(t1)++) + { + return THRUST_FWD(t1)++; + } +}; // end postfix_increment template<typename Eval> __host__ __device__ actor< composite< - unary_operator<suffix_increment>, + transparent_unary_operator<postfix_increment>, actor<Eval> > > operator++(const actor<Eval> &_1, int) { - return compose(unary_operator<suffix_increment>(), _1); + return compose(transparent_unary_operator<postfix_increment>(), _1); } // end operator++() + // there's no standard prefix_decrement functional, so roll an ad hoc one here -template<typename T> - struct prefix_decrement - : public thrust::unary_function<T&,T&> +struct prefix_decrement { - __host__ __device__ T& operator()(T &x) const { return --x; } + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1> + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(--THRUST_FWD(t1))) -> decltype(--THRUST_FWD(t1)) + { + return --THRUST_FWD(t1); + } }; // end prefix_decrement template<typename Eval> __host__ __device__ actor< composite< - unary_operator<prefix_decrement>, + transparent_unary_operator<prefix_decrement>, actor<Eval> > > operator--(const actor<Eval> &_1) { - return compose(unary_operator<prefix_decrement>(), _1); + return compose(transparent_unary_operator<prefix_decrement>(), _1); } // end operator--() -// there's no standard suffix_decrement functional, so roll an ad hoc one here -template<typename T> - struct suffix_decrement - : public thrust::unary_function<T&,T> + +// there's no standard postfix_decrement functional, so roll an ad hoc one here +struct postfix_decrement { - __host__ __device__ T operator()(T &x) const { return x--; } -}; // end suffix_decrement + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1> + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(THRUST_FWD(t1)--)) -> decltype(THRUST_FWD(t1)--) + { + return THRUST_FWD(t1)--; + } +}; // end prefix_increment template<typename Eval> __host__ __device__ actor< composite< - unary_operator<suffix_decrement>, + transparent_unary_operator<postfix_decrement>, actor<Eval> > > operator--(const actor<Eval> &_1, int) { - return compose(unary_operator<suffix_decrement>(), _1); + return compose(transparent_unary_operator<postfix_decrement>(), _1); } // end operator--() } // end functional diff --git a/thrust/detail/functional/operators/assignment_operator.h b/thrust/detail/functional/operators/assignment_operator.h index fb8958f88b..a2f18339bc 100644 --- a/thrust/detail/functional/operators/assignment_operator.h +++ b/thrust/detail/functional/operators/assignment_operator.h @@ -37,19 +37,27 @@ namespace functional template<typename> struct as_actor; // there's no standard assign functional, so roll an ad hoc one here -template<typename T> - struct assign - : thrust::binary_function<T&,T,T&> +struct assign { - __host__ __device__ T& operator()(T &lhs, const T &rhs) const { return lhs = rhs; } -}; // end assign + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) = THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) = THRUST_FWD(t2)) + { + return THRUST_FWD(t1) = THRUST_FWD(t2); + } +}; template<typename Eval, typename T> struct assign_result { typedef actor< composite< - binary_operator<assign>, + transparent_binary_operator<assign>, actor<Eval>, typename as_actor<T>::type > @@ -61,7 +69,7 @@ template<typename Eval, typename T> typename assign_result<Eval,T>::type do_assign(const actor<Eval> &_1, const T &_2) { - return compose(binary_operator<assign>(), + return compose(transparent_binary_operator<assign>(), _1, as_actor<T>::convert(_2)); } // end do_assign() diff --git a/thrust/detail/functional/operators/bitwise_operators.h b/thrust/detail/functional/operators/bitwise_operators.h index 796f1701c1..a6461f9d49 100644 --- a/thrust/detail/functional/operators/bitwise_operators.h +++ b/thrust/detail/functional/operators/bitwise_operators.h @@ -33,14 +33,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_and>, + transparent_binary_operator<bit_and<>>, actor<T1>, typename as_actor<T2>::type > > operator&(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::bit_and>(), + return compose(transparent_binary_operator<bit_and<>>(), make_actor(_1), make_actor(_2)); } // end operator&() @@ -49,14 +49,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_and>, + transparent_binary_operator<bit_and<>>, typename as_actor<T1>::type, actor<T2> > > operator&(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::bit_and>(), + return compose(transparent_binary_operator<bit_and<>>(), make_actor(_1), make_actor(_2)); } // end operator&() @@ -65,14 +65,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_and>, + transparent_binary_operator<bit_and<>>, actor<T1>, actor<T2> > > operator&(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::bit_and>(), + return compose(transparent_binary_operator<bit_and<>>(), make_actor(_1), make_actor(_2)); } // end operator&() @@ -81,14 +81,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_or>, + transparent_binary_operator<bit_or<>>, actor<T1>, typename as_actor<T2>::type > > operator|(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::bit_or>(), + return compose(transparent_binary_operator<bit_or<>>(), make_actor(_1), make_actor(_2)); } // end operator|() @@ -97,14 +97,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_or>, + transparent_binary_operator<bit_or<>>, typename as_actor<T1>::type, actor<T2> > > operator|(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::bit_or>(), + return compose(transparent_binary_operator<bit_or<>>(), make_actor(_1), make_actor(_2)); } // end operator|() @@ -113,14 +113,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_or>, + transparent_binary_operator<bit_or<>>, actor<T1>, actor<T2> > > operator|(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::bit_or>(), + return compose(transparent_binary_operator<bit_or<>>(), make_actor(_1), make_actor(_2)); } // end operator|() @@ -129,14 +129,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_xor>, + transparent_binary_operator<bit_xor<>>, actor<T1>, typename as_actor<T2>::type > > operator^(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::bit_xor>(), + return compose(transparent_binary_operator<bit_xor<>>(), make_actor(_1), make_actor(_2)); } // end operator^() @@ -145,14 +145,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_xor>, + transparent_binary_operator<bit_xor<>>, typename as_actor<T1>::type, actor<T2> > > operator^(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::bit_xor>(), + return compose(transparent_binary_operator<bit_xor<>>(), make_actor(_1), make_actor(_2)); } // end operator^() @@ -161,60 +161,76 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::bit_xor>, + transparent_binary_operator<bit_xor<>>, actor<T1>, actor<T2> > > operator^(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::bit_xor>(), + return compose(transparent_binary_operator<bit_xor<>>(), make_actor(_1), make_actor(_2)); } // end operator^() + // there's no standard bit_not functional, so roll an ad hoc one here -template<typename T> - struct bit_not - : public thrust::unary_function<T,T> +struct bit_not { - __host__ __device__ T operator()(const T &x) const {return ~x;} -}; // end bit_not + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1> + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(~THRUST_FWD(t1))) -> decltype(~THRUST_FWD(t1)) + { + return ~THRUST_FWD(t1); + } +}; // end prefix_increment template<typename Eval> __host__ __device__ actor< composite< - unary_operator<bit_not>, + transparent_unary_operator<bit_not>, actor<Eval> > > __host__ __device__ operator~(const actor<Eval> &_1) { - return compose(unary_operator<bit_not>(), _1); + return compose(transparent_unary_operator<bit_not>(), _1); } // end operator~() // there's no standard bit_lshift functional, so roll an ad hoc one here -template<typename T> - struct bit_lshift - : public thrust::binary_function<T,T,T> +struct bit_lshift { - __host__ __device__ T operator()(const T &lhs, const T &rhs) const {return lhs << rhs;} -}; // end bit_lshift + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) << THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) << THRUST_FWD(t2)) + { + return THRUST_FWD(t1) << THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_lshift>, + transparent_binary_operator<bit_lshift>, actor<T1>, typename as_actor<T2>::type > > operator<<(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_lshift>(), + return compose(transparent_binary_operator<bit_lshift>(), make_actor(_1), make_actor(_2)); } // end operator<<() @@ -223,14 +239,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_lshift>, + transparent_binary_operator<bit_lshift>, typename as_actor<T1>::type, actor<T2> > > operator<<(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_lshift>(), + return compose(transparent_binary_operator<bit_lshift>(), make_actor(_1), make_actor(_2)); } // end operator<<() @@ -239,38 +255,47 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_lshift>, + transparent_binary_operator<bit_lshift>, actor<T1>, actor<T2> > > operator<<(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_lshift>(), + return compose(transparent_binary_operator<bit_lshift>(), make_actor(_1), make_actor(_2)); } // end operator<<() // there's no standard bit_rshift functional, so roll an ad hoc one here -template<typename T> - struct bit_rshift - : public thrust::binary_function<T,T,T> +struct bit_rshift { - __host__ __device__ T operator()(const T &lhs, const T &rhs) const {return lhs >> rhs;} -}; // end bit_rshift + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) >> THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) >> THRUST_FWD(t2)) + { + return THRUST_FWD(t1) >> THRUST_FWD(t2); + } +}; + template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_rshift>, + transparent_binary_operator<bit_rshift>, actor<T1>, typename as_actor<T2>::type > > operator>>(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_rshift>(), + return compose(transparent_binary_operator<bit_rshift>(), make_actor(_1), make_actor(_2)); } // end operator>>() @@ -279,14 +304,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_rshift>, + transparent_binary_operator<bit_rshift>, typename as_actor<T1>::type, actor<T2> > > operator>>(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_rshift>(), + return compose(transparent_binary_operator<bit_rshift>(), make_actor(_1), make_actor(_2)); } // end operator>>() @@ -295,14 +320,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_rshift>, + transparent_binary_operator<bit_rshift>, actor<T1>, actor<T2> > > operator>>(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_rshift>(), + return compose(transparent_binary_operator<bit_rshift>(), make_actor(_1), make_actor(_2)); } // end operator>>() diff --git a/thrust/detail/functional/operators/compound_assignment_operators.h b/thrust/detail/functional/operators/compound_assignment_operators.h index cb8d4c1050..737d6abd09 100644 --- a/thrust/detail/functional/operators/compound_assignment_operators.h +++ b/thrust/detail/functional/operators/compound_assignment_operators.h @@ -28,25 +28,34 @@ namespace detail namespace functional { -template<typename T> - struct plus_equal - : public thrust::binary_function<T&,T,T&> -{ - __host__ __device__ T& operator()(T &lhs, const T &rhs) const { return lhs += rhs; } -}; // end plus_equal +// there's no standard plus_equal functional, so roll an ad hoc one here +struct plus_equal +{ + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) += THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) += THRUST_FWD(t2)) + { + return THRUST_FWD(t1) += THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<plus_equal>, + transparent_binary_operator<plus_equal>, actor<T1>, typename as_actor<T2>::type > > operator+=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<plus_equal>(), + return compose(transparent_binary_operator<plus_equal>(), make_actor(_1), make_actor(_2)); } // end operator+=() @@ -55,37 +64,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<plus_equal>, + transparent_binary_operator<plus_equal>, actor<T1>, actor<T2> > > operator+=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<plus_equal>(), + return compose(transparent_binary_operator<plus_equal>(), make_actor(_1), make_actor(_2)); } // end operator+=() -template<typename T> - struct minus_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard minus_equal functional, so roll an ad hoc one here +struct minus_equal { - __host__ __device__ T& operator()(T &lhs, const T &rhs) const { return lhs -= rhs; } -}; // end minus_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) -= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) -= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) -= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<minus_equal>, + transparent_binary_operator<minus_equal>, actor<T1>, typename as_actor<T2>::type > > operator-=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<minus_equal>(), + return compose(transparent_binary_operator<minus_equal>(), make_actor(_1), make_actor(_2)); } // end operator-=() @@ -94,37 +112,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<minus_equal>, + transparent_binary_operator<minus_equal>, actor<T1>, actor<T2> > > operator-=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<minus_equal>(), + return compose(transparent_binary_operator<minus_equal>(), make_actor(_1), make_actor(_2)); } // end operator-=() -template<typename T> - struct multiplies_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard multiplies_equal functional, so roll an ad hoc one here +struct multiplies_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs *= rhs; } -}; // end multiplies_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) *= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) *= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) *= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<multiplies_equal>, + transparent_binary_operator<multiplies_equal>, actor<T1>, typename as_actor<T2>::type > > operator*=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<multiplies_equal>(), + return compose(transparent_binary_operator<multiplies_equal>(), make_actor(_1), make_actor(_2)); } // end operator*=() @@ -133,37 +160,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<multiplies_equal>, + transparent_binary_operator<multiplies_equal>, actor<T1>, actor<T2> > > operator*=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<multiplies_equal>(), + return compose(transparent_binary_operator<multiplies_equal>(), make_actor(_1), make_actor(_2)); } // end operator*=() -template<typename T> - struct divides_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard divides_equal functional, so roll an ad hoc one here +struct divides_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs /= rhs; } -}; // end divides_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) /= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) /= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) /= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<divides_equal>, + transparent_binary_operator<divides_equal>, actor<T1>, typename as_actor<T2>::type > > operator/=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<divides_equal>(), + return compose(transparent_binary_operator<divides_equal>(), make_actor(_1), make_actor(_2)); } // end operator/=() @@ -172,37 +208,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<divides_equal>, + transparent_binary_operator<divides_equal>, actor<T1>, actor<T2> > > operator/=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<divides_equal>(), + return compose(transparent_binary_operator<divides_equal>(), make_actor(_1), make_actor(_2)); } // end operator/=() -template<typename T> - struct modulus_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard modulus_equal functional, so roll an ad hoc one here +struct modulus_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs %= rhs; } -}; // end modulus_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) %= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) %= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) %= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<modulus_equal>, + transparent_binary_operator<modulus_equal>, actor<T1>, typename as_actor<T2>::type > > operator%=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<modulus_equal>(), + return compose(transparent_binary_operator<modulus_equal>(), make_actor(_1), make_actor(_2)); } // end operator%=() @@ -211,37 +256,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<modulus_equal>, + transparent_binary_operator<modulus_equal>, actor<T1>, actor<T2> > > operator%=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<modulus_equal>(), + return compose(transparent_binary_operator<modulus_equal>(), make_actor(_1), make_actor(_2)); } // end operator%=() -template<typename T> - struct bit_and_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard bit_and_equal functional, so roll an ad hoc one here +struct bit_and_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs &= rhs; } -}; // end bit_and_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) &= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) &= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) &= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_and_equal>, + transparent_binary_operator<bit_and_equal>, actor<T1>, typename as_actor<T2>::type > > operator&=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_and_equal>(), + return compose(transparent_binary_operator<bit_and_equal>(), make_actor(_1), make_actor(_2)); } // end operator&=() @@ -250,37 +304,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_and_equal>, + transparent_binary_operator<bit_and_equal>, actor<T1>, actor<T2> > > operator&=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_and_equal>(), + return compose(transparent_binary_operator<bit_and_equal>(), make_actor(_1), make_actor(_2)); } // end operator&=() -template<typename T> - struct bit_or_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard bit_or_equal functional, so roll an ad hoc one here +struct bit_or_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs |= rhs; } -}; // end bit_or_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) |= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) |= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) |= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_or_equal>, + transparent_binary_operator<bit_or_equal>, actor<T1>, typename as_actor<T2>::type > > operator|=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_or_equal>(), + return compose(transparent_binary_operator<bit_or_equal>(), make_actor(_1), make_actor(_2)); } // end operator|=() @@ -289,37 +352,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_or_equal>, + transparent_binary_operator<bit_or_equal>, actor<T1>, actor<T2> > > operator|=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_or_equal>(), + return compose(transparent_binary_operator<bit_or_equal>(), make_actor(_1), make_actor(_2)); } // end operator|=() -template<typename T> - struct bit_xor_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard bit_xor_equal functional, so roll an ad hoc one here +struct bit_xor_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs ^= rhs; } -}; // end bit_xor_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) ^= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) ^= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) ^= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_xor_equal>, + transparent_binary_operator<bit_xor_equal>, actor<T1>, typename as_actor<T2>::type > > operator^=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_xor_equal>(), + return compose(transparent_binary_operator<bit_xor_equal>(), make_actor(_1), make_actor(_2)); } // end operator|=() @@ -328,37 +400,45 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_xor_equal>, + transparent_binary_operator<bit_xor_equal>, actor<T1>, actor<T2> > > operator^=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_xor_equal>(), + return compose(transparent_binary_operator<bit_xor_equal>(), make_actor(_1), make_actor(_2)); } // end operator|=() -template<typename T> - struct bit_lshift_equal - : public thrust::binary_function<T&,T,T&> -{ - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs <<= rhs; } -}; // end bit_lshift_equal - +// there's no standard bit_lshift_equal functional, so roll an ad hoc one here +struct bit_lshift_equal +{ + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) <<= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) <<= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) <<= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_lshift_equal>, + transparent_binary_operator<bit_lshift_equal>, actor<T1>, typename as_actor<T2>::type > > operator<<=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_lshift_equal>(), + return compose(transparent_binary_operator<bit_lshift_equal>(), make_actor(_1), make_actor(_2)); } // end operator<<=() @@ -367,37 +447,46 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_lshift_equal>, + transparent_binary_operator<bit_lshift_equal>, actor<T1>, actor<T2> > > operator<<=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_lshift_equal>(), + return compose(transparent_binary_operator<bit_lshift_equal>(), make_actor(_1), make_actor(_2)); } // end operator<<=() -template<typename T> - struct bit_rshift_equal - : public thrust::binary_function<T&,T,T&> +// there's no standard bit_rshift_equal functional, so roll an ad hoc one here +struct bit_rshift_equal { - __host__ __device__ T& operator()(T &lhs, const T&rhs) const { return lhs >>= rhs; } -}; // end bit_rshift_equal + using is_transparent = void; + + __thrust_exec_check_disable__ + template <typename T1, typename T2> + __host__ __device__ + constexpr auto operator()(T1&& t1, T2&& t2) const + noexcept(noexcept(THRUST_FWD(t1) >>= THRUST_FWD(t2))) + -> decltype(THRUST_FWD(t1) >>= THRUST_FWD(t2)) + { + return THRUST_FWD(t1) >>= THRUST_FWD(t2); + } +}; template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_rshift_equal>, + transparent_binary_operator<bit_rshift_equal>, actor<T1>, typename as_actor<T2>::type > > operator>>=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<bit_rshift_equal>(), + return compose(transparent_binary_operator<bit_rshift_equal>(), make_actor(_1), make_actor(_2)); } // end operator>>=() @@ -406,14 +495,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<bit_rshift_equal>, + transparent_binary_operator<bit_rshift_equal>, actor<T1>, actor<T2> > > operator>>=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<bit_rshift_equal>(), + return compose(transparent_binary_operator<bit_rshift_equal>(), make_actor(_1), make_actor(_2)); } // end operator>>=() diff --git a/thrust/detail/functional/operators/logical_operators.h b/thrust/detail/functional/operators/logical_operators.h index f5e39e125e..85a2e5e040 100644 --- a/thrust/detail/functional/operators/logical_operators.h +++ b/thrust/detail/functional/operators/logical_operators.h @@ -33,14 +33,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::logical_and>, + transparent_binary_operator<thrust::logical_and<>>, actor<T1>, typename as_actor<T2>::type > > operator&&(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::logical_and>(), + return compose(transparent_binary_operator<thrust::logical_and<>>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -49,14 +49,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::logical_and>, + transparent_binary_operator<thrust::logical_and<>>, typename as_actor<T1>::type, actor<T2> > > operator&&(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::logical_and>(), + return compose(transparent_binary_operator<thrust::logical_and<>>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -65,14 +65,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::logical_and>, + transparent_binary_operator<thrust::logical_and<>>, actor<T1>, actor<T2> > > operator&&(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::logical_and>(), + return compose(transparent_binary_operator<thrust::logical_and<>>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -81,14 +81,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::logical_or>, + transparent_binary_operator<thrust::logical_or<>>, actor<T1>, typename as_actor<T2>::type > > operator||(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::logical_or>(), + return compose(transparent_binary_operator<thrust::logical_or<>>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -97,14 +97,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::logical_or>, + transparent_binary_operator<thrust::logical_or<>>, typename as_actor<T1>::type, actor<T2> > > operator||(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::logical_or>(), + return compose(transparent_binary_operator<thrust::logical_or<>>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -113,14 +113,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::logical_or>, + transparent_binary_operator<thrust::logical_or<>>, actor<T1>, actor<T2> > > operator||(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::logical_or>(), + return compose(transparent_binary_operator<thrust::logical_or<>>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -129,13 +129,13 @@ template<typename Eval> __host__ __device__ actor< composite< - unary_operator<thrust::logical_not>, + transparent_unary_operator<thrust::logical_not<>>, actor<Eval> > > operator!(const actor<Eval> &_1) { - return compose(unary_operator<thrust::logical_not>(), _1); + return compose(transparent_unary_operator<thrust::logical_not<>>(), _1); } // end operator!() } // end functional diff --git a/thrust/detail/functional/operators/operator_adaptors.h b/thrust/detail/functional/operators/operator_adaptors.h index 6649211130..873e68b61d 100644 --- a/thrust/detail/functional/operators/operator_adaptors.h +++ b/thrust/detail/functional/operators/operator_adaptors.h @@ -17,8 +17,13 @@ #pragma once #include <thrust/detail/config.h> +#include <thrust/detail/functional/argument.h> +#include <thrust/detail/type_deduction.h> #include <thrust/tuple.h> #include <thrust/detail/type_traits.h> +#include <thrust/type_traits/void_t.h> + +#include <type_traits> namespace thrust { @@ -27,87 +32,104 @@ namespace detail namespace functional { -// this thing (which models Eval) is an adaptor for the unary -// functors inside functional.h -template<template<typename> class UnaryOperator> - struct unary_operator +// Adapts a transparent unary functor from functional.h (e.g. thrust::negate<>) +// into the Eval interface. +template <typename UnaryFunctor> +struct transparent_unary_operator { - template<typename Env> - struct argument - : thrust::detail::eval_if< - (thrust::tuple_size<Env>::value == 0), - thrust::detail::identity_<thrust::null_type>, - thrust::tuple_element<0,Env> - > - { - }; + template <typename> + using operator_type = UnaryFunctor; + + template <typename Env> + using argument = + typename thrust::detail::eval_if< + thrust::tuple_size<Env>::value != 1, + thrust::detail::identity_<thrust::null_type>, + thrust::detail::functional::argument_helper<0, Env> + >::type; - template<typename Env> - struct operator_type + template <typename Env> + struct result_type_impl { - typedef UnaryOperator< - typename thrust::detail::remove_reference< - typename argument<Env>::type - >::type - > type; + using type = decltype( + std::declval<UnaryFunctor>()(std::declval<argument<Env>>())); }; - template<typename Env> - struct result + template <typename Env> + using result_type = + typename thrust::detail::eval_if< + std::is_same<thrust::null_type, argument<Env>>::value, + thrust::detail::identity_<thrust::null_type>, + result_type_impl<Env> + >::type; + + template <typename Env> + struct result { - typedef typename operator_type<Env>::type op_type; - typedef typename op_type::result_type type; + using op_type = UnaryFunctor; + using type = result_type<Env>; }; - template<typename Env> + template <typename Env> __host__ __device__ - typename result<Env>::type eval(const Env &e) const - { - typename operator_type<Env>::type op; - return op(thrust::get<0>(e)); - } // end eval() -}; // end unary_operator - -// this thing (which models Eval) is an adaptor for the binary -// functors inside functional.h -template<template<typename> class BinaryOperator> - struct binary_operator + result_type<Env> eval(Env&& e) const + THRUST_RETURNS(UnaryFunctor{}(thrust::get<0>(THRUST_FWD(e)))) +}; + + +// Adapts a transparent unary functor from functional.h (e.g. thrust::less<>) +// into the Eval interface. +template <typename BinaryFunctor> +struct transparent_binary_operator { - template<typename Env> - struct first_argument - : thrust::detail::eval_if< - (thrust::tuple_size<Env>::value == 0), - thrust::detail::identity_<thrust::null_type>, - thrust::tuple_element<0,Env> - > - { - }; + template <typename> + using operator_type = BinaryFunctor; + + template <typename Env> + using first_argument = + typename thrust::detail::eval_if< + thrust::tuple_size<Env>::value != 2, + thrust::detail::identity_<thrust::null_type>, + thrust::detail::functional::argument_helper<0, Env> + >::type; - template<typename Env> - struct operator_type + template <typename Env> + using second_argument = + typename thrust::detail::eval_if< + thrust::tuple_size<Env>::value != 2, + thrust::detail::identity_<thrust::null_type>, + thrust::detail::functional::argument_helper<1, Env> + >::type; + + template <typename Env> + struct result_type_impl { - typedef BinaryOperator< - typename thrust::detail::remove_reference< - typename first_argument<Env>::type - >::type - > type; + using type = decltype( + std::declval<BinaryFunctor>()(std::declval<first_argument<Env>>(), + std::declval<second_argument<Env>>())); }; - template<typename Env> - struct result + template <typename Env> + using result_type = + typename thrust::detail::eval_if< + (std::is_same<thrust::null_type, first_argument<Env>>::value || + std::is_same<thrust::null_type, second_argument<Env>>::value), + thrust::detail::identity_<thrust::null_type>, + result_type_impl<Env> + >::type; + + template <typename Env> + struct result { - typedef typename operator_type<Env>::type op_type; - typedef typename op_type::result_type type; + using op_type = BinaryFunctor; + using type = result_type<Env>; }; - template<typename Env> + template <typename Env> __host__ __device__ - typename result<Env>::type eval(const Env &e) const - { - typename operator_type<Env>::type op; - return op(thrust::get<0>(e), thrust::get<1>(e)); - } // end eval() -}; // end binary_operator + result_type<Env> eval(Env&& e) const + THRUST_RETURNS(BinaryFunctor{}(thrust::get<0>(e), thrust::get<1>(e))) +}; } // end functional } // end detail diff --git a/thrust/detail/functional/operators/relational_operators.h b/thrust/detail/functional/operators/relational_operators.h index ec8864715b..51fd4640a2 100644 --- a/thrust/detail/functional/operators/relational_operators.h +++ b/thrust/detail/functional/operators/relational_operators.h @@ -33,14 +33,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::equal_to>, + transparent_binary_operator<thrust::equal_to<>>, actor<T1>, typename as_actor<T2>::type > > operator==(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::equal_to>(), + return compose(transparent_binary_operator<thrust::equal_to<>>(), make_actor(_1), make_actor(_2)); } // end operator==() @@ -49,14 +49,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::equal_to>, + transparent_binary_operator<thrust::equal_to<>>, typename as_actor<T1>::type, actor<T2> > > operator==(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::equal_to>(), + return compose(transparent_binary_operator<thrust::equal_to<>>(), make_actor(_1), make_actor(_2)); } // end operator==() @@ -65,14 +65,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::equal_to>, + transparent_binary_operator<thrust::equal_to<>>, actor<T1>, actor<T2> > > operator==(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::equal_to>(), + return compose(transparent_binary_operator<thrust::equal_to<>>(), make_actor(_1), make_actor(_2)); } // end operator==() @@ -81,14 +81,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::not_equal_to>, + transparent_binary_operator<thrust::not_equal_to<>>, actor<T1>, typename as_actor<T2>::type > > operator!=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::not_equal_to>(), + return compose(transparent_binary_operator<thrust::not_equal_to<>>(), make_actor(_1), make_actor(_2)); } // end operator!=() @@ -97,14 +97,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::not_equal_to>, + transparent_binary_operator<thrust::not_equal_to<>>, typename as_actor<T1>::type, actor<T2> > > operator!=(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::not_equal_to>(), + return compose(transparent_binary_operator<thrust::not_equal_to<>>(), make_actor(_1), make_actor(_2)); } // end operator!=() @@ -113,14 +113,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::not_equal_to>, + transparent_binary_operator<thrust::not_equal_to<>>, actor<T1>, actor<T2> > > operator!=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::not_equal_to>(), + return compose(transparent_binary_operator<thrust::not_equal_to<>>(), make_actor(_1), make_actor(_2)); } // end operator!=() @@ -129,14 +129,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::greater>, + transparent_binary_operator<thrust::greater<>>, actor<T1>, typename as_actor<T2>::type > > operator>(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::greater>(), + return compose(transparent_binary_operator<thrust::greater<>>(), make_actor(_1), make_actor(_2)); } // end operator>() @@ -145,14 +145,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::greater>, + transparent_binary_operator<thrust::greater<>>, typename as_actor<T1>::type, actor<T2> > > operator>(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::greater>(), + return compose(transparent_binary_operator<thrust::greater<>>(), make_actor(_1), make_actor(_2)); } // end operator>() @@ -161,14 +161,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::greater>, + transparent_binary_operator<thrust::greater<>>, actor<T1>, actor<T2> > > operator>(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::greater>(), + return compose(transparent_binary_operator<thrust::greater<>>(), make_actor(_1), make_actor(_2)); } // end operator>() @@ -177,14 +177,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::less>, + transparent_binary_operator<thrust::less<>>, actor<T1>, typename as_actor<T2>::type > > operator<(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::less>(), + return compose(transparent_binary_operator<thrust::less<>>(), make_actor(_1), make_actor(_2)); } // end operator<() @@ -193,14 +193,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::less>, + transparent_binary_operator<thrust::less<>>, typename as_actor<T1>::type, actor<T2> > > operator<(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::less>(), + return compose(transparent_binary_operator<thrust::less<>>(), make_actor(_1), make_actor(_2)); } // end operator<() @@ -209,14 +209,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::less>, + transparent_binary_operator<thrust::less<>>, actor<T1>, actor<T2> > > operator<(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::less>(), + return compose(transparent_binary_operator<thrust::less<>>(), make_actor(_1), make_actor(_2)); } // end operator<() @@ -225,14 +225,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::greater_equal>, + transparent_binary_operator<thrust::greater_equal<>>, actor<T1>, typename as_actor<T2>::type > > operator>=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::greater_equal>(), + return compose(transparent_binary_operator<thrust::greater_equal<>>(), make_actor(_1), make_actor(_2)); } // end operator>=() @@ -241,14 +241,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::greater_equal>, + transparent_binary_operator<thrust::greater_equal<>>, typename as_actor<T1>::type, actor<T2> > > operator>=(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::greater_equal>(), + return compose(transparent_binary_operator<thrust::greater_equal<>>(), make_actor(_1), make_actor(_2)); } // end operator>=() @@ -257,14 +257,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::greater_equal>, + transparent_binary_operator<thrust::greater_equal<>>, actor<T1>, actor<T2> > > operator>=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::greater_equal>(), + return compose(transparent_binary_operator<thrust::greater_equal<>>(), make_actor(_1), make_actor(_2)); } // end operator>=() @@ -273,14 +273,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::less_equal>, + transparent_binary_operator<thrust::less_equal<>>, actor<T1>, typename as_actor<T2>::type > > operator<=(const actor<T1> &_1, const T2 &_2) { - return compose(binary_operator<thrust::less_equal>(), + return compose(transparent_binary_operator<thrust::less_equal<>>(), make_actor(_1), make_actor(_2)); } // end operator<=() @@ -289,14 +289,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::less_equal>, + transparent_binary_operator<thrust::less_equal<>>, typename as_actor<T1>::type, actor<T2> > > operator<=(const T1 &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::less_equal>(), + return compose(transparent_binary_operator<thrust::less_equal<>>(), make_actor(_1), make_actor(_2)); } // end operator<=() @@ -305,14 +305,14 @@ template<typename T1, typename T2> __host__ __device__ actor< composite< - binary_operator<thrust::less_equal>, + transparent_binary_operator<thrust::less_equal<>>, actor<T1>, actor<T2> > > operator<=(const actor<T1> &_1, const actor<T2> &_2) { - return compose(binary_operator<thrust::less_equal>(), + return compose(transparent_binary_operator<thrust::less_equal<>>(), make_actor(_1), make_actor(_2)); } // end operator<=()