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