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
serial_adjacent_find(InputIterator first,InputIterator last,Compare compare,command_queue & queue)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
adjacent_find_with_atomics(InputIterator first,InputIterator last,Compare compare,command_queue & queue)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
adjacent_find(InputIterator first,InputIterator last,Compare compare,command_queue & queue=system::default_queue ())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
adjacent_find(InputIterator first,InputIterator last,command_queue & queue=system::default_queue ())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
170