From a2451e308f29c60c81b90b5ec86b59087464e2d0 Mon Sep 17 00:00:00 2001 From: Allison Vacanti 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 | 60 +--- 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, 604 insertions(+), 501 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 +#include #include #include @@ -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 v(1000); + thrust::sequence(v.begin(), v.end()); + thrust::device_vector 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 #include + +#include #include #include #include +#include template 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 v1(100, 1.f); + thrust::device_vector v2(100, 1.f); + + auto result = thrust::inner_product(v1.begin(), v1.end(), v2.begin(), 0.0f, + thrust::plus{}, + _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..01e8d5cd35 100644 --- a/thrust/detail/functional/actor.h +++ b/thrust/detail/functional/actor.h @@ -30,6 +30,7 @@ #include #include #include +#include #include namespace thrust @@ -39,6 +40,14 @@ namespace detail namespace functional { +// eval_ref is +// - T when T is a subclass of thrust::reference +// - T& otherwise +// This is used to let thrust::references pass through actor evaluations. +template +using eval_ref = typename std::conditional< + thrust::detail::is_wrapped_reference::value, T, T&>::type; + template struct apply_actor { @@ -61,55 +70,10 @@ template typename apply_actor::type operator()(void) const; - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7) const; - - template - __host__ __device__ - typename apply_actor >::type - operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8) const; - - template + template __host__ __device__ - typename apply_actor >::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...>>::type + operator()(Ts&&... ts) const; template __host__ __device__ diff --git a/thrust/detail/functional/actor.inl b/thrust/detail/functional/actor.inl index 2c7fadd366..444d2ff1a5 100644 --- a/thrust/detail/functional/actor.inl +++ b/thrust/detail/functional/actor.inl @@ -27,6 +27,9 @@ #include #include #include +#include + +#include namespace thrust { @@ -62,135 +65,38 @@ template return eval_type::eval(thrust::null_type()); } // end basic_environment::operator() -template - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::operator()(T0 &_0) const -{ - return eval_type::eval(thrust::tie(_0)); -} // end basic_environment::operator() +// actor::operator() needs to construct a tuple of references to its +// arguments. To make this work with thrust::reference, 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, or +// - T& for any other types. +// This struct provides a nicer diagnostic for when these conditions aren't +// met. +template +using actor_check_ref_type = + thrust::detail::integral_constant::value || + thrust::detail::is_wrapped_reference::value )>; + +template +using actor_check_ref_types = + thrust::conjunction...>; template - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::operator()(T0 &_0, T1 &_1) const +template +__host__ __device__ +typename apply_actor::eval_type, + thrust::tuple...>>::type +actor::operator()(Ts&&... ts) const { - return eval_type::eval(thrust::tie(_0,_1)); -} // end basic_environment::operator() - -template - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::operator()(T0 &_0, T1 &_1, T2 &_2) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2)); -} // end basic_environment::operator() - -template - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::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 - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::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 - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::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 - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::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 - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::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 - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5,_6,_7,_8)); -} // end basic_environment::operator() - -template - template - __host__ __device__ - typename apply_actor< - typename actor::eval_type, - typename thrust::tuple - >::type - actor - ::operator()(T0 &_0, T1 &_1, T2 &_2, T3 &_3, T4 &_4, T5 &_5, T6 &_6, T7 &_7, T8 &_8, T9 &_9) const -{ - return eval_type::eval(thrust::tie(_0,_1,_2,_3,_4,_5,_6,_7,_8,_9)); -} // end basic_environment::operator() + static_assert(actor_check_ref_types::value, + "Actor evaluations only support rvalue references to " + "thrust::reference subclasses."); + using tuple_type = thrust::tuple...>; + return eval_type::eval(tuple_type(THRUST_FWD(ts)...)); +} // end actor::operator() template template 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 __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator>, actor > > __host__ __device__ operator-(const actor &_1) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator>(), _1); } // end operator-() // there's no standard unary_plus functional, so roll an ad hoc one here -template - struct unary_plus - : public thrust::unary_function +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 + __host__ __device__ + constexpr auto operator()(T1&& t1) const + noexcept(noexcept(+THRUST_FWD(t1))) -> decltype(+THRUST_FWD(t1)) + { + return +THRUST_FWD(t1); + } +}; template __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator, actor > > operator+(const actor &_1) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator(), _1); } // end operator+() template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator+(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator+() @@ -84,14 +91,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator+(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator+() @@ -100,14 +107,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator+(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator+() @@ -116,14 +123,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator-(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator-() @@ -132,14 +139,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator-(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator-() @@ -148,14 +155,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator-(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator-() @@ -164,14 +171,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator*(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator*() @@ -180,14 +187,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator*(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator*() @@ -196,14 +203,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator*(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator*() @@ -212,14 +219,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator/(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator/() @@ -228,14 +235,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator/(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator/() @@ -244,14 +251,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator/(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator/() @@ -260,14 +267,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator%(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator%() @@ -276,14 +283,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator%(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator%() @@ -292,100 +299,131 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator%(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator%() // there's no standard prefix_increment functional, so roll an ad hoc one here -template - struct prefix_increment - : public thrust::unary_function +struct prefix_increment { - __host__ __device__ T& operator()(T &x) const { return ++x; } + using is_transparent = void; + + __thrust_exec_check_disable__ + template + __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 __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator, actor > > operator++(const actor &_1) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator(), _1); } // end operator++() -// there's no standard suffix_increment functional, so roll an ad hoc one here -template - struct suffix_increment - : public thrust::unary_function + +// 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 + __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 __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator, actor > > operator++(const actor &_1, int) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator(), _1); } // end operator++() + // there's no standard prefix_decrement functional, so roll an ad hoc one here -template - struct prefix_decrement - : public thrust::unary_function +struct prefix_decrement { - __host__ __device__ T& operator()(T &x) const { return --x; } + using is_transparent = void; + + __thrust_exec_check_disable__ + template + __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 __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator, actor > > operator--(const actor &_1) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator(), _1); } // end operator--() -// there's no standard suffix_decrement functional, so roll an ad hoc one here -template - struct suffix_decrement - : public thrust::unary_function + +// 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 + __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 __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator, actor > > operator--(const actor &_1, int) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator(), _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 struct as_actor; // there's no standard assign functional, so roll an ad hoc one here -template - struct assign - : thrust::binary_function +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 + __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 struct assign_result { typedef actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > @@ -61,7 +69,7 @@ template typename assign_result::type do_assign(const actor &_1, const T &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), _1, as_actor::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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator&(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&() @@ -49,14 +49,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator&(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&() @@ -65,14 +65,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator&(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&() @@ -81,14 +81,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator|(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator|() @@ -97,14 +97,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator|(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator|() @@ -113,14 +113,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator|(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator|() @@ -129,14 +129,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator^(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator^() @@ -145,14 +145,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator^(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator^() @@ -161,60 +161,76 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator^(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator^() + // there's no standard bit_not functional, so roll an ad hoc one here -template - struct bit_not - : public thrust::unary_function +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 + __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 __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator, actor > > __host__ __device__ operator~(const actor &_1) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator(), _1); } // end operator~() // there's no standard bit_lshift functional, so roll an ad hoc one here -template - struct bit_lshift - : public thrust::binary_function +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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator<<(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator<<() @@ -223,14 +239,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, typename as_actor::type, actor > > operator<<(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator<<() @@ -239,38 +255,47 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator<<(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator<<() // there's no standard bit_rshift functional, so roll an ad hoc one here -template - struct bit_rshift - : public thrust::binary_function +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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator>>(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator>>() @@ -279,14 +304,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, typename as_actor::type, actor > > operator>>(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator>>() @@ -295,14 +320,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator>>(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), 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 - struct plus_equal - : public thrust::binary_function -{ - __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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator+=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator+=() @@ -55,37 +64,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator+=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator+=() -template - struct minus_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator-=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator-=() @@ -94,37 +112,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator-=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator-=() -template - struct multiplies_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator*=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator*=() @@ -133,37 +160,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator*=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator*=() -template - struct divides_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator/=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator/=() @@ -172,37 +208,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator/=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator/=() -template - struct modulus_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator%=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator%=() @@ -211,37 +256,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator%=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator%=() -template - struct bit_and_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator&=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator&=() @@ -250,37 +304,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator&=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator&=() -template - struct bit_or_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator|=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator|=() @@ -289,37 +352,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator|=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator|=() -template - struct bit_xor_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator^=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator|=() @@ -328,37 +400,45 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator^=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator|=() -template - struct bit_lshift_equal - : public thrust::binary_function -{ - __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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator<<=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator<<=() @@ -367,37 +447,46 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator<<=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator<<=() -template - struct bit_rshift_equal - : public thrust::binary_function +// 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 + __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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, typename as_actor::type > > operator>>=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), make_actor(_1), make_actor(_2)); } // end operator>>=() @@ -406,14 +495,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator, actor, actor > > operator>>=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator(), 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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator&&(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -49,14 +49,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator&&(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -65,14 +65,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator&&(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -81,14 +81,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator||(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -97,14 +97,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator||(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -113,14 +113,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator||(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator&&() @@ -129,13 +129,13 @@ template __host__ __device__ actor< composite< - unary_operator, + transparent_unary_operator>, actor > > operator!(const actor &_1) { - return compose(unary_operator(), _1); + return compose(transparent_unary_operator>(), _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 +#include +#include #include #include +#include + +#include 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 class UnaryOperator> - struct unary_operator +// Adapts a transparent unary functor from functional.h (e.g. thrust::negate<>) +// into the Eval interface. +template +struct transparent_unary_operator { - template - struct argument - : thrust::detail::eval_if< - (thrust::tuple_size::value == 0), - thrust::detail::identity_, - thrust::tuple_element<0,Env> - > - { - }; + template + using operator_type = UnaryFunctor; + + template + using argument = + typename thrust::detail::eval_if< + thrust::tuple_size::value != 1, + thrust::detail::identity_, + thrust::detail::functional::argument_helper<0, Env> + >::type; - template - struct operator_type + template + struct result_type_impl { - typedef UnaryOperator< - typename thrust::detail::remove_reference< - typename argument::type - >::type - > type; + using type = decltype( + std::declval()(std::declval>())); }; - template - struct result + template + using result_type = + typename thrust::detail::eval_if< + std::is_same>::value, + thrust::detail::identity_, + result_type_impl + >::type; + + template + struct result { - typedef typename operator_type::type op_type; - typedef typename op_type::result_type type; + using op_type = UnaryFunctor; + using type = result_type; }; - template + template __host__ __device__ - typename result::type eval(const Env &e) const - { - typename operator_type::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 class BinaryOperator> - struct binary_operator + result_type 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 +struct transparent_binary_operator { - template - struct first_argument - : thrust::detail::eval_if< - (thrust::tuple_size::value == 0), - thrust::detail::identity_, - thrust::tuple_element<0,Env> - > - { - }; + template + using operator_type = BinaryFunctor; + + template + using first_argument = + typename thrust::detail::eval_if< + thrust::tuple_size::value != 2, + thrust::detail::identity_, + thrust::detail::functional::argument_helper<0, Env> + >::type; - template - struct operator_type + template + using second_argument = + typename thrust::detail::eval_if< + thrust::tuple_size::value != 2, + thrust::detail::identity_, + thrust::detail::functional::argument_helper<1, Env> + >::type; + + template + struct result_type_impl { - typedef BinaryOperator< - typename thrust::detail::remove_reference< - typename first_argument::type - >::type - > type; + using type = decltype( + std::declval()(std::declval>(), + std::declval>())); }; - template - struct result + template + using result_type = + typename thrust::detail::eval_if< + (std::is_same>::value || + std::is_same>::value), + thrust::detail::identity_, + result_type_impl + >::type; + + template + struct result { - typedef typename operator_type::type op_type; - typedef typename op_type::result_type type; + using op_type = BinaryFunctor; + using type = result_type; }; - template + template __host__ __device__ - typename result::type eval(const Env &e) const - { - typename operator_type::type op; - return op(thrust::get<0>(e), thrust::get<1>(e)); - } // end eval() -}; // end binary_operator + result_type 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 __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator==(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator==() @@ -49,14 +49,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator==(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator==() @@ -65,14 +65,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator==(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator==() @@ -81,14 +81,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator!=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator!=() @@ -97,14 +97,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator!=(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator!=() @@ -113,14 +113,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator!=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator!=() @@ -129,14 +129,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator>(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator>() @@ -145,14 +145,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator>(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator>() @@ -161,14 +161,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator>(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator>() @@ -177,14 +177,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator<(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator<() @@ -193,14 +193,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator<(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator<() @@ -209,14 +209,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator<(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator<() @@ -225,14 +225,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator>=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator>=() @@ -241,14 +241,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator>=(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator>=() @@ -257,14 +257,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator>=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator>=() @@ -273,14 +273,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, typename as_actor::type > > operator<=(const actor &_1, const T2 &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator<=() @@ -289,14 +289,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, typename as_actor::type, actor > > operator<=(const T1 &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator<=() @@ -305,14 +305,14 @@ template __host__ __device__ actor< composite< - binary_operator, + transparent_binary_operator>, actor, actor > > operator<=(const actor &_1, const actor &_2) { - return compose(binary_operator(), + return compose(transparent_binary_operator>(), make_actor(_1), make_actor(_2)); } // end operator<=()