#pragma once #include #if THRUST_CPP_DIALECT >= 2014 #include #include #include namespace testing { namespace async { namespace exclusive_scan { namespace mixin { //------------------------------------------------------------------------------ namespace postfix_args { template > struct all_overloads { using postfix_args_type = std::tuple< // List any extra arg overloads: std::tuple<>, // - no extra args std::tuple, // - initial_value std::tuple // - initial_value, binary_op >; static postfix_args_type generate_postfix_args() { return postfix_args_type{std::tuple<>{}, std::make_tuple(value_type{42}), std::make_tuple(value_type{42}, alternate_binary_op{})}; } }; } // namespace postfix_args //------------------------------------------------------------------------------ namespace invoke_reference { template struct host_synchronous { template static void invoke_reference(InputType const& input, OutputType& output, PostfixArgTuple&& postfix_tuple, std::index_sequence) { // Create host versions of the input/output: thrust::host_vector host_input(input.cbegin(), input.cend()); thrust::host_vector host_output(host_input.size()); // Run host synchronous algorithm to generate reference. thrust::exclusive_scan(host_input.cbegin(), host_input.cend(), host_output.begin(), std::get( THRUST_FWD(postfix_tuple))...); // Copy back to device. output = host_output; } }; } // namespace invoke_reference //------------------------------------------------------------------------------ namespace invoke_async { struct simple { template static auto invoke_async(PrefixArgTuple&& prefix_tuple, std::index_sequence, InputType const& input, OutputType& output, PostfixArgTuple&& postfix_tuple, std::index_sequence) { auto e = thrust::async::exclusive_scan( std::get(THRUST_FWD(prefix_tuple))..., input.cbegin(), input.cend(), output.begin(), std::get(THRUST_FWD(postfix_tuple))...); return e; } }; } // namespace invoke_async } // namespace mixin } // namespace exclusive_scan } // namespace async } // namespace testing #endif // C++14