1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2014 Roshan <thisisroshansmail@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_SET_SYMMETRIC_DIFFERENCE_HPP
12 #define BOOST_COMPUTE_ALGORITHM_SET_SYMMETRIC_DIFFERENCE_HPP
13
14 #include <iterator>
15
16 #include <boost/static_assert.hpp>
17
18 #include <boost/compute/algorithm/detail/compact.hpp>
19 #include <boost/compute/algorithm/detail/balanced_path.hpp>
20 #include <boost/compute/algorithm/exclusive_scan.hpp>
21 #include <boost/compute/algorithm/fill_n.hpp>
22 #include <boost/compute/container/vector.hpp>
23 #include <boost/compute/detail/iterator_range_size.hpp>
24 #include <boost/compute/detail/meta_kernel.hpp>
25 #include <boost/compute/system.hpp>
26 #include <boost/compute/type_traits/is_device_iterator.hpp>
27
28 namespace boost {
29 namespace compute {
30 namespace detail {
31
32 ///
33 /// \brief Serial set symmetric difference kernel class
34 ///
35 /// Subclass of meta_kernel to perform serial set symmetric
36 /// difference after tiling
37 ///
38 class serial_set_symmetric_difference_kernel : meta_kernel
39 {
40 public:
41 unsigned int tile_size;
42
serial_set_symmetric_difference_kernel()43 serial_set_symmetric_difference_kernel() : meta_kernel("set_symmetric_difference")
44 {
45 tile_size = 4;
46 }
47
48 template<class InputIterator1, class InputIterator2,
49 class InputIterator3, class InputIterator4,
50 class OutputIterator1, class OutputIterator2>
set_range(InputIterator1 first1,InputIterator2 first2,InputIterator3 tile_first1,InputIterator3 tile_last1,InputIterator4 tile_first2,OutputIterator1 result,OutputIterator2 counts)51 void set_range(InputIterator1 first1,
52 InputIterator2 first2,
53 InputIterator3 tile_first1,
54 InputIterator3 tile_last1,
55 InputIterator4 tile_first2,
56 OutputIterator1 result,
57 OutputIterator2 counts)
58 {
59 m_count = iterator_range_size(tile_first1, tile_last1) - 1;
60
61 *this <<
62 "uint i = get_global_id(0);\n" <<
63 "uint start1 = " << tile_first1[expr<uint_>("i")] << ";\n" <<
64 "uint end1 = " << tile_first1[expr<uint_>("i+1")] << ";\n" <<
65 "uint start2 = " << tile_first2[expr<uint_>("i")] << ";\n" <<
66 "uint end2 = " << tile_first2[expr<uint_>("i+1")] << ";\n" <<
67 "uint index = i*" << tile_size << ";\n" <<
68 "uint count = 0;\n" <<
69 "while(start1<end1 && start2<end2)\n" <<
70 "{\n" <<
71 " if(" << first1[expr<uint_>("start1")] << " == " <<
72 first2[expr<uint_>("start2")] << ")\n" <<
73 " {\n" <<
74 " start1++; start2++;\n" <<
75 " }\n" <<
76 " else if(" << first1[expr<uint_>("start1")] << " < " <<
77 first2[expr<uint_>("start2")] << ")\n" <<
78 " {\n" <<
79 result[expr<uint_>("index")] <<
80 " = " << first1[expr<uint_>("start1")] << ";\n" <<
81 " index++; count++;\n" <<
82 " start1++;\n" <<
83 " }\n" <<
84 " else\n" <<
85 " {\n" <<
86 result[expr<uint_>("index")] <<
87 " = " << first2[expr<uint_>("start2")] << ";\n" <<
88 " index++; count++;\n" <<
89 " start2++;\n" <<
90 " }\n" <<
91 "}\n" <<
92 "while(start1<end1)\n" <<
93 "{\n" <<
94 result[expr<uint_>("index")] <<
95 " = " << first1[expr<uint_>("start1")] << ";\n" <<
96 " index++; count++;\n" <<
97 " start1++;\n" <<
98 "}\n" <<
99 "while(start2<end2)\n" <<
100 "{\n" <<
101 result[expr<uint_>("index")] <<
102 " = " << first2[expr<uint_>("start2")] << ";\n" <<
103 " index++; count++;\n" <<
104 " start2++;\n" <<
105 "}\n" <<
106 counts[expr<uint_>("i")] << " = count;\n";
107 }
108
exec(command_queue & queue)109 event exec(command_queue &queue)
110 {
111 if(m_count == 0) {
112 return event();
113 }
114
115 return exec_1d(queue, 0, m_count);
116 }
117
118 private:
119 size_t m_count;
120 };
121
122 } //end detail namespace
123
124 ///
125 /// \brief Set symmetric difference algorithm
126 ///
127 /// Finds the symmetric difference of the sorted range [first2, last2) from
128 /// the sorted range [first1, last1) and stores it in range starting at result
129 /// \return Iterator pointing to end of symmetric difference
130 ///
131 /// \param first1 Iterator pointing to start of first set
132 /// \param last1 Iterator pointing to end of first set
133 /// \param first2 Iterator pointing to start of second set
134 /// \param last2 Iterator pointing to end of second set
135 /// \param result Iterator pointing to start of range in which the symmetric
136 /// difference will be stored
137 /// \param queue Queue on which to execute
138 ///
139 /// Space complexity:
140 /// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
141 template<class InputIterator1, class InputIterator2, class OutputIterator>
142 inline OutputIterator
set_symmetric_difference(InputIterator1 first1,InputIterator1 last1,InputIterator2 first2,InputIterator2 last2,OutputIterator result,command_queue & queue=system::default_queue ())143 set_symmetric_difference(InputIterator1 first1,
144 InputIterator1 last1,
145 InputIterator2 first2,
146 InputIterator2 last2,
147 OutputIterator result,
148 command_queue &queue = system::default_queue())
149 {
150 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
151 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
152 BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
153
154 typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
155
156 int tile_size = 1024;
157
158 int count1 = detail::iterator_range_size(first1, last1);
159 int count2 = detail::iterator_range_size(first2, last2);
160
161 vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
162 vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
163
164 // Tile the sets
165 detail::balanced_path_kernel tiling_kernel;
166 tiling_kernel.tile_size = tile_size;
167 tiling_kernel.set_range(first1, last1, first2, last2,
168 tile_a.begin()+1, tile_b.begin()+1);
169 fill_n(tile_a.begin(), 1, 0, queue);
170 fill_n(tile_b.begin(), 1, 0, queue);
171 tiling_kernel.exec(queue);
172
173 fill_n(tile_a.end()-1, 1, count1, queue);
174 fill_n(tile_b.end()-1, 1, count2, queue);
175
176 vector<value_type> temp_result(count1+count2, queue.get_context());
177 vector<uint_> counts((count1+count2+tile_size-1)/tile_size + 1, queue.get_context());
178 fill_n(counts.end()-1, 1, 0, queue);
179
180 // Find individual symmetric differences
181 detail::serial_set_symmetric_difference_kernel symmetric_difference_kernel;
182 symmetric_difference_kernel.tile_size = tile_size;
183 symmetric_difference_kernel.set_range(first1, first2, tile_a.begin(),
184 tile_a.end(), tile_b.begin(),
185 temp_result.begin(), counts.begin());
186
187 symmetric_difference_kernel.exec(queue);
188
189 exclusive_scan(counts.begin(), counts.end(), counts.begin(), queue);
190
191 // Compact the results
192 detail::compact_kernel compact_kernel;
193 compact_kernel.tile_size = tile_size;
194 compact_kernel.set_range(temp_result.begin(), counts.begin(), counts.end(), result);
195
196 compact_kernel.exec(queue);
197
198 return result + (counts.end() - 1).read(queue);
199 }
200
201 } //end compute namespace
202 } //end boost namespace
203
204 #endif // BOOST_COMPUTE_ALGORITHM_SET_SYMMETRIC_DIFFERENCE_HPP
205