compute/algorithm/adjacent_find.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@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_ADJACENT_FIND_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_ADJACENT_FIND_HPP
0013
0014 #include <iterator>
0015
0016 #include <boost/static_assert.hpp>
0017
0018 #include <boost/compute/command_queue.hpp>
0019 #include <boost/compute/lambda.hpp>
0020 #include <boost/compute/system.hpp>
0021 #include <boost/compute/container/detail/scalar.hpp>
0022 #include <boost/compute/detail/iterator_range_size.hpp>
0023 #include <boost/compute/detail/meta_kernel.hpp>
0024 #include <boost/compute/functional/operator.hpp>
0025 #include <boost/compute/type_traits/vector_size.hpp>
0026 #include <boost/compute/type_traits/is_device_iterator.hpp>
0027
0028 namespace boost {
0029 namespace compute {
0030 namespace detail {
0031
0032 template<class InputIterator, class Compare>
0033 inline InputIterator
0034 serial_adjacent_find(InputIterator first,
0035                      InputIterator last,
0036                      Compare compare,
0037                      command_queue &queue)
0038 {
0039     if(first == last){
0040         return last;
0041     }
0042
0043     const context &context = queue.get_context();
0044
0045     detail::scalar<uint_> output(context);
0046
0047     detail::meta_kernel k("serial_adjacent_find");
0048
0049     size_t size_arg = k.add_arg<const uint_>("size");
0050     size_t output_arg = k.add_arg<uint_ *>(memory_object::global_memory, "output");
0051
0052     k << k.decl<uint_>("result") << " = size;\n"
0053       << "for(uint i = 0; i < size - 1; i++){\n"
0054       << "    if(" << compare(first[k.expr<uint_>("i")],
0055                               first[k.expr<uint_>("i+1")]) << "){\n"
0056       << "        result = i;\n"
0057       << "        break;\n"
0058       << "    }\n"
0059       << "}\n"
0060       << "*output = result;\n";
0061
0062     k.set_arg<const uint_>(
0063         size_arg, static_cast<uint_>(detail::iterator_range_size(first, last))
0064     );
0065     k.set_arg(output_arg, output.get_buffer());
0066
0067     k.exec_1d(queue, 0, 1, 1);
0068
0069     return first + output.read(queue);
0070 }
0071
0072 template<class InputIterator, class Compare>
0073 inline InputIterator
0074 adjacent_find_with_atomics(InputIterator first,
0075                            InputIterator last,
0076                            Compare compare,
0077                            command_queue &queue)
0078 {
0079     if(first == last){
0080         return last;
0081     }
0082
0083     const context &context = queue.get_context();
0084     size_t count = detail::iterator_range_size(first, last);
0085
0086     // initialize output to the last index
0087     detail::scalar<uint_> output(context);
0088     output.write(static_cast<uint_>(count), queue);
0089
0090     detail::meta_kernel k("adjacent_find_with_atomics");
0091
0092     size_t output_arg = k.add_arg<uint_ *>(memory_object::global_memory, "output");
0093
0094     k << "const uint i = get_global_id(0);\n"
0095       << "if(" << compare(first[k.expr<uint_>("i")],
0096                           first[k.expr<uint_>("i+1")]) << "){\n"
0097       << "    atomic_min(output, i);\n"
0098       << "}\n";
0099
0100     k.set_arg(output_arg, output.get_buffer());
0101
0102     k.exec_1d(queue, 0, count - 1, 1);
0103
0104     return first + output.read(queue);
0105 }
0106
0107 } // end detail namespace
0108
0109 /// Searches the range [\p first, \p last) for two identical adjacent
0110 /// elements and returns an iterator pointing to the first.
0111 ///
0112 /// \param first first element in the range to search
0113 /// \param last last element in the range to search
0114 /// \param compare binary comparison function
0115 /// \param queue command queue to perform the operation
0116 ///
0117 /// \return \c InputIteratorm to the first element which compares equal
0118 ///         to the following element. If none are equal, returns \c last.
0119 ///
0120 /// Space complexity: \Omega(1)
0121 ///
0122 /// \see find(), adjacent_difference()
0123 template<class InputIterator, class Compare>
0124 inline InputIterator
0125 adjacent_find(InputIterator first,
0126               InputIterator last,
0127               Compare compare,
0128               command_queue &queue = system::default_queue())
0129 {
0130     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0131     size_t count = detail::iterator_range_size(first, last);
0132     if(count < 32){
0133         return detail::serial_adjacent_find(first, last, compare, queue);
0134     }
0135     else {
0136         return detail::adjacent_find_with_atomics(first, last, compare, queue);
0137     }
0138 }
0139
0140 /// \overload
0141 template<class InputIterator>
0142 inline InputIterator
0143 adjacent_find(InputIterator first,
0144               InputIterator last,
0145               command_queue &queue = system::default_queue())
0146 {
0147     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0148     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0149
0150     using ::boost::compute::lambda::_1;
0151     using ::boost::compute::lambda::_2;
0152     using ::boost::compute::lambda::all;
0153
0154     if(vector_size<value_type>::value == 1){
0155         return ::boost::compute::adjacent_find(
0156             first, last, _1 == _2, queue
0157         );
0158     }
0159     else {
0160         return ::boost::compute::adjacent_find(
0161             first, last, all(_1 == _2), queue
0162         );
0163     }
0164 }
0165
0166 } // end compute namespace
0167 } // end boost namespace
0168
0169 #endif // BOOST_COMPUTE_ALGORITHM_ADJACENT_FIND_HPP