/* * Copyright 2008-2021 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. * You may obtain a copy of the License at * * http://www.apache.org/licenses/LICENSE-2.0 * * Unless required by applicable law or agreed to in writing, software * distributed under the License is distributed on an "AS IS" BASIS, * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. * See the License for the specific language governing permissions and * limitations under the License. */ /*! \file * \brief Algorithms for asynchronously reducing a range to a single value. */ #pragma once #include #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) # pragma GCC system_header #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) # pragma clang system_header #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header #include #if _CCCL_STD_VER >= 2014 # include # include # include # include # include # include # include THRUST_NAMESPACE_BEGIN namespace async { /*! \cond */ namespace unimplemented { template _CCCL_HOST future async_reduce(thrust::execution_policy&, ForwardIt, Sentinel, T, BinaryOp) { THRUST_STATIC_ASSERT_MSG((thrust::detail::depend_on_instantiation::value), "this algorithm is not implemented for the specified system"); return {}; } } // namespace unimplemented namespace reduce_detail { using thrust::async::unimplemented::async_reduce; struct reduce_fn final { template _CCCL_HOST static auto call(thrust::detail::execution_policy_base const& exec, ForwardIt&& first, Sentinel&& last, T&& init, BinaryOp&& op) // ADL dispatch. THRUST_RETURNS(async_reduce( thrust::detail::derived_cast(thrust::detail::strip_const(exec)), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(init), THRUST_FWD(op))) template _CCCL_HOST static auto call4( thrust::detail::execution_policy_base const& exec, ForwardIt&& first, Sentinel&& last, T&& init, thrust::true_type) // ADL dispatch. THRUST_RETURNS(async_reduce( thrust::detail::derived_cast(thrust::detail::strip_const(exec)), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(init), thrust::plus>{})) template _CCCL_HOST static auto call3(thrust::detail::execution_policy_base const& exec, ForwardIt&& first, Sentinel&& last, thrust::true_type) // ADL dispatch. THRUST_RETURNS(async_reduce( thrust::detail::derived_cast(thrust::detail::strip_const(exec)), THRUST_FWD(first), THRUST_FWD(last), typename iterator_traits>::value_type{}, thrust::plus>::value_type>>{})) template _CCCL_HOST static auto call4(ForwardIt&& first, Sentinel&& last, T&& init, BinaryOp&& op, thrust::false_type) THRUST_RETURNS(reduce_fn::call( thrust::detail::select_system(typename iterator_system>::type{}), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(init), THRUST_FWD(op))) template _CCCL_HOST static auto call3(ForwardIt&& first, Sentinel&& last, T&& init, thrust::false_type) THRUST_RETURNS(reduce_fn::call( thrust::detail::select_system(typename iterator_system>::type{}), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(init), thrust::plus>{})) // MSVC WAR: MSVC gets angsty and eats all available RAM when we try to detect // if T1 is an execution_policy by using SFINAE. Switching to a static // dispatch pattern to prevent this. template _CCCL_HOST static auto call(T1&& t1, T2&& t2, T3&& t3) THRUST_RETURNS(reduce_fn::call3( THRUST_FWD(t1), THRUST_FWD(t2), THRUST_FWD(t3), thrust::is_execution_policy>{})) template _CCCL_HOST static auto call(T1&& t1, T2&& t2, T3&& t3, T4&& t4) THRUST_RETURNS(reduce_fn::call4( THRUST_FWD(t1), THRUST_FWD(t2), THRUST_FWD(t3), THRUST_FWD(t4), thrust::is_execution_policy>{})) template _CCCL_HOST static auto call(ForwardIt&& first, Sentinel&& last) THRUST_RETURNS(reduce_fn::call( thrust::detail::select_system(typename iterator_system>::type{}), THRUST_FWD(first), THRUST_FWD(last), typename iterator_traits>::value_type{}, thrust::plus>::value_type>>{})) template _CCCL_NODISCARD _CCCL_HOST auto operator()(Args&&... args) const THRUST_RETURNS(call(THRUST_FWD(args)...)) }; } // namespace reduce_detail THRUST_INLINE_CONSTANT reduce_detail::reduce_fn reduce{}; /////////////////////////////////////////////////////////////////////////////// namespace unimplemented { template _CCCL_HOST event async_reduce_into(thrust::execution_policy&, ForwardIt, Sentinel, OutputIt, T, BinaryOp) { THRUST_STATIC_ASSERT_MSG((thrust::detail::depend_on_instantiation::value), "this algorithm is not implemented for the specified system"); return {}; } } // namespace unimplemented namespace reduce_into_detail { using thrust::async::unimplemented::async_reduce_into; struct reduce_into_fn final { template _CCCL_HOST static auto call(thrust::detail::execution_policy_base const& exec, ForwardIt&& first, Sentinel&& last, OutputIt&& output, T&& init, BinaryOp&& op) // ADL dispatch. THRUST_RETURNS(async_reduce_into( thrust::detail::derived_cast(thrust::detail::strip_const(exec)), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), THRUST_FWD(init), THRUST_FWD(op))) template _CCCL_HOST static auto call5( thrust::detail::execution_policy_base const& exec, ForwardIt&& first, Sentinel&& last, OutputIt&& output, T&& init, thrust::true_type) // ADL dispatch. THRUST_RETURNS(async_reduce_into( thrust::detail::derived_cast(thrust::detail::strip_const(exec)), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), THRUST_FWD(init), thrust::plus>{})) template _CCCL_HOST static auto call4( thrust::detail::execution_policy_base const& exec, ForwardIt&& first, Sentinel&& last, OutputIt&& output, thrust::true_type) // ADL dispatch. THRUST_RETURNS(async_reduce_into( thrust::detail::derived_cast(thrust::detail::strip_const(exec)), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), typename iterator_traits>::value_type{}, thrust::plus>::value_type>>{})) template _CCCL_HOST static auto call5( ForwardIt&& first, Sentinel&& last, OutputIt&& output, T&& init, BinaryOp&& op, thrust::false_type) THRUST_RETURNS(reduce_into_fn::call( thrust::detail::select_system(typename iterator_system>::type{}, typename iterator_system>::type{}), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), THRUST_FWD(init), THRUST_FWD(op))) template _CCCL_HOST static auto call4( ForwardIt&& first, Sentinel&& last, OutputIt&& output, T&& init, thrust::false_type) THRUST_RETURNS(reduce_into_fn::call( thrust::detail::select_system(typename iterator_system>::type{}, typename iterator_system>::type{}), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), THRUST_FWD(init), thrust::plus>{})) template _CCCL_HOST static auto call(ForwardIt&& first, Sentinel&& last, OutputIt&& output) THRUST_RETURNS(reduce_into_fn::call( thrust::detail::select_system(typename iterator_system>::type{}, typename iterator_system>::type{}), THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), typename iterator_traits>::value_type{}, thrust::plus>::value_type>>{})) // MSVC WAR: MSVC gets angsty and eats all available RAM when we try to detect // if T1 is an execution_policy by using SFINAE. Switching to a static // dispatch pattern to prevent this. template _CCCL_HOST static auto call(T1&& t1, T2&& t2, T3&& t3, T4&& t4) THRUST_RETURNS(reduce_into_fn::call4( THRUST_FWD(t1), THRUST_FWD(t2), THRUST_FWD(t3), THRUST_FWD(t4), thrust::is_execution_policy>{})) template _CCCL_HOST static auto call(T1&& t1, T2&& t2, T3&& t3, T4&& t4, T5&& t5) THRUST_RETURNS(reduce_into_fn::call5( THRUST_FWD(t1), THRUST_FWD(t2), THRUST_FWD(t3), THRUST_FWD(t4), THRUST_FWD(t5), thrust::is_execution_policy>{})) template _CCCL_NODISCARD _CCCL_HOST auto operator()(Args&&... args) const THRUST_RETURNS(call(THRUST_FWD(args)...)) }; } // namespace reduce_into_detail THRUST_INLINE_CONSTANT reduce_into_detail::reduce_into_fn reduce_into{}; /*! \endcond */ } // namespace async THRUST_NAMESPACE_END #endif