algorithm/detail/reduce_on_cpu.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2016 Jakub Szuppe <j.szuppe@gmail.com>
0003 //
0004 // Distributed under the Boost Software License, Version 1.0
0005 // See accompanying file LICENSE_1_0.txt or copy at
0006 // http://www.boost.org/LICENSE_1_0.txt
0007 //
0008 // See http://boostorg.github.com/compute for more information.
0009 //---------------------------------------------------------------------------//
0010
0011 #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP
0013
0014 #include <algorithm>
0015
0016 #include <boost/compute/buffer.hpp>
0017 #include <boost/compute/command_queue.hpp>
0018 #include <boost/compute/detail/meta_kernel.hpp>
0019 #include <boost/compute/detail/iterator_range_size.hpp>
0020 #include <boost/compute/detail/parameter_cache.hpp>
0021 #include <boost/compute/iterator/buffer_iterator.hpp>
0022 #include <boost/compute/type_traits/result_of.hpp>
0023 #include <boost/compute/algorithm/detail/serial_reduce.hpp>
0024
0025 namespace boost {
0026 namespace compute {
0027 namespace detail {
0028
0029 template<class InputIterator, class OutputIterator, class BinaryFunction>
0030 inline void reduce_on_cpu(InputIterator first,
0031                           InputIterator last,
0032                           OutputIterator result,
0033                           BinaryFunction function,
0034                           command_queue &queue)
0035 {
0036     typedef typename
0037         std::iterator_traits<InputIterator>::value_type T;
0038     typedef typename
0039         ::boost::compute::result_of<BinaryFunction(T, T)>::type result_type;
0040
0041     const device &device = queue.get_device();
0042     const uint_ compute_units = queue.get_device().compute_units();
0043
0044     boost::shared_ptr<parameter_cache> parameters =
0045         detail::parameter_cache::get_global_cache(device);
0046
0047     std::string cache_key =
0048         "__boost_reduce_cpu_" + boost::lexical_cast<std::string>(sizeof(T));
0049
0050     // for inputs smaller than serial_reduce_threshold
0051     // serial_reduce algorithm is used
0052     uint_ serial_reduce_threshold =
0053         parameters->get(cache_key, "serial_reduce_threshold", 16384 * sizeof(T));
0054     serial_reduce_threshold =
0055         (std::max)(serial_reduce_threshold, uint_(compute_units));
0056
0057     const context &context = queue.get_context();
0058     size_t count = detail::iterator_range_size(first, last);
0059     if(count == 0){
0060         return;
0061     }
0062     else if(count < serial_reduce_threshold) {
0063         return serial_reduce(first, last, result, function, queue);
0064     }
0065
0066     meta_kernel k("reduce_on_cpu");
0067     buffer output(context, sizeof(result_type) * compute_units);
0068
0069     size_t count_arg = k.add_arg<uint_>("count");
0070     size_t output_arg =
0071         k.add_arg<result_type *>(memory_object::global_memory, "output");
0072
0073     k <<
0074         "uint block = " <<
0075             "(uint)ceil(((float)count)/get_global_size(0));\n" <<
0076         "uint index = get_global_id(0) * block;\n" <<
0077         "uint end = min(count, index + block);\n" <<
0078
0079         k.decl<result_type>("result") << " = " << first[k.var<uint_>("index")] << ";\n" <<
0080         "index++;\n" <<
0081         "while(index < end){\n" <<
0082              "result = " << function(k.var<T>("result"),
0083                                      first[k.var<uint_>("index")]) << ";\n" <<
0084              "index++;\n" <<
0085         "}\n" <<
0086         "output[get_global_id(0)] = result;\n";
0087
0088     size_t global_work_size = compute_units;
0089     kernel kernel = k.compile(context);
0090
0091     // reduction to global_work_size elements
0092     kernel.set_arg(count_arg, static_cast<uint_>(count));
0093     kernel.set_arg(output_arg, output);
0094     queue.enqueue_1d_range_kernel(kernel, 0, global_work_size, 0);
0095
0096     // final reduction
0097     reduce_on_cpu(
0098         make_buffer_iterator<result_type>(output),
0099         make_buffer_iterator<result_type>(output, global_work_size),
0100         result,
0101         function,
0102         queue
0103     );
0104 }
0105
0106 } // end detail namespace
0107 } // end compute namespace
0108 } // end boost namespace
0109
0110 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP