algorithm/detail/find_extrema_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_FIND_EXTREMA_ON_CPU_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_ON_CPU_HPP
0013
0014 #include <algorithm>
0015
0016 #include <boost/compute/algorithm/detail/find_extrema_with_reduce.hpp>
0017 #include <boost/compute/algorithm/detail/find_extrema_with_atomics.hpp>
0018 #include <boost/compute/algorithm/detail/serial_find_extrema.hpp>
0019 #include <boost/compute/detail/iterator_range_size.hpp>
0020 #include <boost/compute/iterator/buffer_iterator.hpp>
0021
0022 namespace boost {
0023 namespace compute {
0024 namespace detail {
0025
0026 template<class InputIterator, class Compare>
0027 inline InputIterator find_extrema_on_cpu(InputIterator first,
0028                                          InputIterator last,
0029                                          Compare compare,
0030                                          const bool find_minimum,
0031                                          command_queue &queue)
0032 {
0033     typedef typename std::iterator_traits<InputIterator>::value_type input_type;
0034     typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
0035     size_t count = iterator_range_size(first, last);
0036
0037     const device &device = queue.get_device();
0038     const uint_ compute_units = queue.get_device().compute_units();
0039
0040     boost::shared_ptr<parameter_cache> parameters =
0041         detail::parameter_cache::get_global_cache(device);
0042     std::string cache_key =
0043         "__boost_find_extrema_cpu_"
0044             + boost::lexical_cast<std::string>(sizeof(input_type));
0045
0046     // for inputs smaller than serial_find_extrema_threshold
0047     // serial_find_extrema algorithm is used
0048     uint_ serial_find_extrema_threshold = parameters->get(
0049         cache_key,
0050         "serial_find_extrema_threshold",
0051         16384 * sizeof(input_type)
0052     );
0053     serial_find_extrema_threshold =
0054         (std::max)(serial_find_extrema_threshold, uint_(2 * compute_units));
0055
0056     const context &context = queue.get_context();
0057     if(count < serial_find_extrema_threshold) {
0058         return serial_find_extrema(first, last, compare, find_minimum, queue);
0059     }
0060
0061     meta_kernel k("find_extrema_on_cpu");
0062     buffer output(context, sizeof(input_type) * compute_units);
0063     buffer output_idx(
0064         context, sizeof(uint_) * compute_units,
0065         buffer::read_write | buffer::alloc_host_ptr
0066     );
0067
0068     size_t count_arg = k.add_arg<uint_>("count");
0069     size_t output_arg =
0070         k.add_arg<input_type *>(memory_object::global_memory, "output");
0071     size_t output_idx_arg =
0072         k.add_arg<uint_ *>(memory_object::global_memory, "output_idx");
0073
0074     k <<
0075         "uint block = " <<
0076             "(uint)ceil(((float)count)/get_global_size(0));\n" <<
0077         "uint index = get_global_id(0) * block;\n" <<
0078         "uint end = min(count, index + block);\n" <<
0079
0080         "uint value_index = index;\n" <<
0081         k.decl<input_type>("value") << " = " << first[k.var<uint_>("index")] << ";\n" <<
0082
0083         "index++;\n" <<
0084         "while(index < end){\n" <<
0085             k.decl<input_type>("candidate") <<
0086                 " = " << first[k.var<uint_>("index")] << ";\n" <<
0087         "#ifndef BOOST_COMPUTE_FIND_MAXIMUM\n" <<
0088             "bool compare = " << compare(k.var<input_type>("candidate"),
0089                                          k.var<input_type>("value")) << ";\n" <<
0090         "#else\n" <<
0091             "bool compare = " << compare(k.var<input_type>("value"),
0092                                          k.var<input_type>("candidate")) << ";\n" <<
0093         "#endif\n" <<
0094             "value = compare ? candidate : value;\n" <<
0095             "value_index = compare ? index : value_index;\n" <<
0096             "index++;\n" <<
0097         "}\n" <<
0098         "output[get_global_id(0)] = value;\n" <<
0099         "output_idx[get_global_id(0)] = value_index;\n";
0100
0101     size_t global_work_size = compute_units;
0102     std::string options;
0103     if(!find_minimum){
0104         options = "-DBOOST_COMPUTE_FIND_MAXIMUM";
0105     }
0106     kernel kernel = k.compile(context, options);
0107
0108     kernel.set_arg(count_arg, static_cast<uint_>(count));
0109     kernel.set_arg(output_arg, output);
0110     kernel.set_arg(output_idx_arg, output_idx);
0111     queue.enqueue_1d_range_kernel(kernel, 0, global_work_size, 0);
0112
0113     buffer_iterator<input_type> result = serial_find_extrema(
0114         make_buffer_iterator<input_type>(output),
0115         make_buffer_iterator<input_type>(output, global_work_size),
0116         compare,
0117         find_minimum,
0118         queue
0119     );
0120
0121     uint_* output_idx_host_ptr =
0122         static_cast<uint_*>(
0123             queue.enqueue_map_buffer(
0124                 output_idx, command_queue::map_read,
0125                 0, global_work_size * sizeof(uint_)
0126             )
0127         );
0128
0129     difference_type extremum_idx =
0130         static_cast<difference_type>(*(output_idx_host_ptr + result.get_index()));
0131     return first + extremum_idx;
0132 }
0133
0134 } // end detail namespace
0135 } // end compute namespace
0136 } // end boost namespace
0137
0138 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_ON_CPU_HPP