]>
Commit | Line | Data |
---|---|---|
7c673cae FG |
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 | /// | |
b32b8144 FG |
117 | /// Space complexity: \Omega(1) |
118 | /// | |
7c673cae FG |
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 |