ceph/src/boost/boost/compute/algorithm/adjacent_find.hpp

   1 //---------------------------------------------------------------------------//
   2 // Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@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_ADJACENT_FIND_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_ADJACENT_FIND_HPP
  13
  14 #include <iterator>
  15
  16 #include <boost/compute/command_queue.hpp>
  17 #include <boost/compute/lambda.hpp>
  18 #include <boost/compute/system.hpp>
  19 #include <boost/compute/container/detail/scalar.hpp>
  20 #include <boost/compute/detail/iterator_range_size.hpp>
  21 #include <boost/compute/detail/meta_kernel.hpp>
  22 #include <boost/compute/functional/operator.hpp>
  23 #include <boost/compute/type_traits/vector_size.hpp>
  24
  25 namespace boost {
  26 namespace compute {
  27 namespace detail {
  28
  29 template<class InputIterator, class Compare>
  30 inline InputIterator
  31 serial_adjacent_find(InputIterator first,
  32                      InputIterator last,
  33                      Compare compare,
  34                      command_queue &queue)
  35 {
  36     if(first == last){
  37         return last;
  38     }
  39
  40     const context &context = queue.get_context();
  41
  42     detail::scalar<uint_> output(context);
  43
  44     detail::meta_kernel k("serial_adjacent_find");
  45
  46     size_t size_arg = k.add_arg<const uint_>("size");
  47     size_t output_arg = k.add_arg<uint_ *>(memory_object::global_memory, "output");
  48
  49     k << k.decl<uint_>("result") << " = size;\n"
  50       << "for(uint i = 0; i < size - 1; i++){\n"
  51       << "    if(" << compare(first[k.expr<uint_>("i")],
  52                               first[k.expr<uint_>("i+1")]) << "){\n"
  53       << "        result = i;\n"
  54       << "        break;\n"
  55       << "    }\n"
  56       << "}\n"
  57       << "*output = result;\n";
  58
  59     k.set_arg<const uint_>(
  60         size_arg, static_cast<uint_>(detail::iterator_range_size(first, last))
  61     );
  62     k.set_arg(output_arg, output.get_buffer());
  63
  64     k.exec_1d(queue, 0, 1, 1);
  65
  66     return first + output.read(queue);
  67 }
  68
  69 template<class InputIterator, class Compare>
  70 inline InputIterator
  71 adjacent_find_with_atomics(InputIterator first,
  72                            InputIterator last,
  73                            Compare compare,
  74                            command_queue &queue)
  75 {
  76     if(first == last){
  77         return last;
  78     }
  79
  80     const context &context = queue.get_context();
  81     size_t count = detail::iterator_range_size(first, last);
  82
  83     // initialize output to the last index
  84     detail::scalar<uint_> output(context);
  85     output.write(static_cast<uint_>(count), queue);
  86
  87     detail::meta_kernel k("adjacent_find_with_atomics");
  88
  89     size_t output_arg = k.add_arg<uint_ *>(memory_object::global_memory, "output");
  90
  91     k << "const uint i = get_global_id(0);\n"
  92       << "if(" << compare(first[k.expr<uint_>("i")],
  93                           first[k.expr<uint_>("i+1")]) << "){\n"
  94       << "    atomic_min(output, i);\n"
  95       << "}\n";
  96
  97     k.set_arg(output_arg, output.get_buffer());
  98
  99     k.exec_1d(queue, 0, count - 1, 1);
 100
 101     return first + output.read(queue);
 102 }
 103
 104 } // end detail namespace
 105
 106 /// Searches the range [\p first, \p last) for two identical adjacent
 107 /// elements and returns an iterator pointing to the first.
 108 ///
 109 /// \param first first element in the range to search
 110 /// \param last last element in the range to search
 111 /// \param compare binary comparison function
 112 /// \param queue command queue to perform the operation
 113 ///
 114 /// \return \c InputIteratorm to the first element which compares equal
 115 ///         to the following element. If none are equal, returns \c last.
 116 ///
 117 /// Space complexity: \Omega(1)
 118 ///
 119 /// \see find(), adjacent_difference()
 120 template<class InputIterator, class Compare>
 121 inline InputIterator
 122 adjacent_find(InputIterator first,
 123               InputIterator last,
 124               Compare compare,
 125               command_queue &queue = system::default_queue())
 126 {
 127     size_t count = detail::iterator_range_size(first, last);
 128     if(count < 32){
 129         return detail::serial_adjacent_find(first, last, compare, queue);
 130     }
 131     else {
 132         return detail::adjacent_find_with_atomics(first, last, compare, queue);
 133     }
 134 }
 135
 136 /// \overload
 137 template<class InputIterator>
 138 inline InputIterator
 139 adjacent_find(InputIterator first,
 140               InputIterator last,
 141               command_queue &queue = system::default_queue())
 142 {
 143     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
 144
 145     using ::boost::compute::lambda::_1;
 146     using ::boost::compute::lambda::_2;
 147     using ::boost::compute::lambda::all;
 148
 149     if(vector_size<value_type>::value == 1){
 150         return ::boost::compute::adjacent_find(
 151             first, last, _1 == _2, queue
 152         );
 153     }
 154     else {
 155         return ::boost::compute::adjacent_find(
 156             first, last, all(_1 == _2), queue
 157         );
 158     }
 159 }
 160
 161 } // end compute namespace
 162 } // end boost namespace
 163
 164 #endif // BOOST_COMPUTE_ALGORITHM_ADJACENT_FIND_HPP