ceph/src/boost/libs/compute/include/boost/compute/algorithm/detail/reduce_on_cpu.hpp

   1 //---------------------------------------------------------------------------//
   2 // Copyright (c) 2016 Jakub Szuppe <j.szuppe@gmail.com>
   3 //
   4 // Distributed under the Boost Software License, Version 1.0
   5 // See accompanying file LICENSE_1_0.txt or copy at
   6 // http://www.boost.org/LICENSE_1_0.txt
   7 //
   8 // See http://boostorg.github.com/compute for more information.
   9 //---------------------------------------------------------------------------//
  10
  11 #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP
  13
  14 #include <algorithm>
  15
  16 #include <boost/compute/buffer.hpp>
  17 #include <boost/compute/command_queue.hpp>
  18 #include <boost/compute/detail/meta_kernel.hpp>
  19 #include <boost/compute/detail/iterator_range_size.hpp>
  20 #include <boost/compute/detail/parameter_cache.hpp>
  21 #include <boost/compute/iterator/buffer_iterator.hpp>
  22 #include <boost/compute/type_traits/result_of.hpp>
  23 #include <boost/compute/algorithm/detail/serial_reduce.hpp>
  24
  25 namespace boost {
  26 namespace compute {
  27 namespace detail {
  28
  29 template<class InputIterator, class OutputIterator, class BinaryFunction>
  30 inline void reduce_on_cpu(InputIterator first,
  31                           InputIterator last,
  32                           OutputIterator result,
  33                           BinaryFunction function,
  34                           command_queue &queue)
  35 {
  36     typedef typename
  37         std::iterator_traits<InputIterator>::value_type T;
  38     typedef typename
  39         ::boost::compute::result_of<BinaryFunction(T, T)>::type result_type;
  40
  41     const device &device = queue.get_device();
  42     const uint_ compute_units = queue.get_device().compute_units();
  43
  44     boost::shared_ptr<parameter_cache> parameters =
  45         detail::parameter_cache::get_global_cache(device);
  46
  47     std::string cache_key =
  48         "__boost_reduce_cpu_" + boost::lexical_cast<std::string>(sizeof(T));
  49
  50     // for inputs smaller than serial_reduce_threshold
  51     // serial_reduce algorithm is used
  52     uint_ serial_reduce_threshold =
  53         parameters->get(cache_key, "serial_reduce_threshold", 16384 * sizeof(T));
  54     serial_reduce_threshold =
  55         (std::max)(serial_reduce_threshold, uint_(compute_units));
  56
  57     const context &context = queue.get_context();
  58     size_t count = detail::iterator_range_size(first, last);
  59     if(count == 0){
  60         return;
  61     }
  62     else if(count < serial_reduce_threshold) {
  63         return serial_reduce(first, last, result, function, queue);
  64     }
  65
  66     meta_kernel k("reduce_on_cpu");
  67     buffer output(context, sizeof(result_type) * compute_units);
  68
  69     size_t count_arg = k.add_arg<uint_>("count");
  70     size_t output_arg =
  71         k.add_arg<result_type *>(memory_object::global_memory, "output");
  72
  73     k <<
  74         "uint block = " <<
  75             "(uint)ceil(((float)count)/get_global_size(0));\n" <<
  76         "uint index = get_global_id(0) * block;\n" <<
  77         "uint end = min(count, index + block);\n" <<
  78
  79         k.decl<result_type>("result") << " = " << first[k.var<uint_>("index")] << ";\n" <<
  80         "index++;\n" <<
  81         "while(index < end){\n" <<
  82              "result = " << function(k.var<T>("result"),
  83                                      first[k.var<uint_>("index")]) << ";\n" <<
  84              "index++;\n" <<
  85         "}\n" <<
  86         "output[get_global_id(0)] = result;\n";
  87
  88     size_t global_work_size = compute_units;
  89     kernel kernel = k.compile(context);
  90
  91     // reduction to global_work_size elements
  92     kernel.set_arg(count_arg, static_cast<uint_>(count));
  93     kernel.set_arg(output_arg, output);
  94     queue.enqueue_1d_range_kernel(kernel, 0, global_work_size, 0);
  95
  96     // final reduction
  97     reduce_on_cpu(
  98         make_buffer_iterator<result_type>(output),
  99         make_buffer_iterator<result_type>(output, global_work_size),
 100         result,
 101         function,
 102         queue
 103     );
 104 }
 105
 106 } // end detail namespace
 107 } // end compute namespace
 108 } // end boost namespace
 109
 110 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP