1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2014 Roshan <thisisroshansmail@gmail.com>
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
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
11 #ifndef BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP
12 #define BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP
16 #include <boost/compute/algorithm/detail/balanced_path.hpp>
17 #include <boost/compute/algorithm/detail/compact.hpp>
18 #include <boost/compute/algorithm/exclusive_scan.hpp>
19 #include <boost/compute/algorithm/fill_n.hpp>
20 #include <boost/compute/container/vector.hpp>
21 #include <boost/compute/detail/iterator_range_size.hpp>
22 #include <boost/compute/detail/meta_kernel.hpp>
23 #include <boost/compute/system.hpp>
30 /// \brief Serial set union kernel class
32 /// Subclass of meta_kernel to perform serial set union after tiling
34 class serial_set_union_kernel : meta_kernel
37 unsigned int tile_size;
39 serial_set_union_kernel() : meta_kernel("set_union")
44 template<class InputIterator1, class InputIterator2,
45 class InputIterator3, class InputIterator4,
46 class OutputIterator1, class OutputIterator2>
47 void set_range(InputIterator1 first1,
48 InputIterator2 first2,
49 InputIterator3 tile_first1,
50 InputIterator3 tile_last1,
51 InputIterator4 tile_first2,
52 OutputIterator1 result,
53 OutputIterator2 counts)
55 m_count = iterator_range_size(tile_first1, tile_last1) - 1;
58 "uint i = get_global_id(0);\n" <<
59 "uint start1 = " << tile_first1[expr<uint_>("i")] << ";\n" <<
60 "uint end1 = " << tile_first1[expr<uint_>("i+1")] << ";\n" <<
61 "uint start2 = " << tile_first2[expr<uint_>("i")] << ";\n" <<
62 "uint end2 = " << tile_first2[expr<uint_>("i+1")] << ";\n" <<
63 "uint index = i*" << tile_size << ";\n" <<
64 "uint count = 0;\n" <<
65 "while(start1<end1 && start2<end2)\n" <<
67 " if(" << first1[expr<uint_>("start1")] << " == " <<
68 first2[expr<uint_>("start2")] << ")\n" <<
70 result[expr<uint_>("index")] <<
71 " = " << first1[expr<uint_>("start1")] << ";\n" <<
72 " index++; count++;\n" <<
73 " start1++; start2++;\n" <<
75 " else if(" << first1[expr<uint_>("start1")] << " < " <<
76 first2[expr<uint_>("start2")] << ")\n" <<
78 result[expr<uint_>("index")] <<
79 " = " << first1[expr<uint_>("start1")] << ";\n" <<
80 " index++; count++;\n" <<
85 result[expr<uint_>("index")] <<
86 " = " << first2[expr<uint_>("start2")] << ";\n" <<
87 " index++; count++;\n" <<
91 "while(start1<end1)\n" <<
93 result[expr<uint_>("index")] <<
94 " = " << first1[expr<uint_>("start1")] << ";\n" <<
95 " index++; count++;\n" <<
98 "while(start2<end2)\n" <<
100 result[expr<uint_>("index")] <<
101 " = " << first2[expr<uint_>("start2")] << ";\n" <<
102 " index++; count++;\n" <<
105 counts[expr<uint_>("i")] << " = count;\n";
108 event exec(command_queue &queue)
114 return exec_1d(queue, 0, m_count);
121 } //end detail namespace
124 /// \brief Set union algorithm
126 /// Finds the union of the sorted range [first1, last1) with the sorted
127 /// range [first2, last2) and stores it in range starting at result
128 /// \return Iterator pointing to end of union
130 /// \param first1 Iterator pointing to start of first set
131 /// \param last1 Iterator pointing to end of first set
132 /// \param first2 Iterator pointing to start of second set
133 /// \param last2 Iterator pointing to end of second set
134 /// \param result Iterator pointing to start of range in which the union
136 /// \param queue Queue on which to execute
138 template<class InputIterator1, class InputIterator2, class OutputIterator>
139 inline OutputIterator set_union(InputIterator1 first1,
140 InputIterator1 last1,
141 InputIterator2 first2,
142 InputIterator2 last2,
143 OutputIterator result,
144 command_queue &queue = system::default_queue())
146 typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
148 int tile_size = 1024;
150 int count1 = detail::iterator_range_size(first1, last1);
151 int count2 = detail::iterator_range_size(first2, last2);
153 vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
154 vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
157 detail::balanced_path_kernel tiling_kernel;
158 tiling_kernel.tile_size = tile_size;
159 tiling_kernel.set_range(first1, last1, first2, last2,
160 tile_a.begin()+1, tile_b.begin()+1);
161 fill_n(tile_a.begin(), 1, 0, queue);
162 fill_n(tile_b.begin(), 1, 0, queue);
163 tiling_kernel.exec(queue);
165 fill_n(tile_a.end()-1, 1, count1, queue);
166 fill_n(tile_b.end()-1, 1, count2, queue);
168 vector<value_type> temp_result(count1+count2, queue.get_context());
169 vector<uint_> counts((count1+count2+tile_size-1)/tile_size + 1, queue.get_context());
170 fill_n(counts.end()-1, 1, 0, queue);
172 // Find individual unions
173 detail::serial_set_union_kernel union_kernel;
174 union_kernel.tile_size = tile_size;
175 union_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(),
176 tile_b.begin(), temp_result.begin(), counts.begin());
178 union_kernel.exec(queue);
180 exclusive_scan(counts.begin(), counts.end(), counts.begin(), queue);
182 // Compact the results
183 detail::compact_kernel compact_kernel;
184 compact_kernel.tile_size = tile_size;
185 compact_kernel.set_range(temp_result.begin(), counts.begin(), counts.end(), result);
187 compact_kernel.exec(queue);
189 return result + (counts.end() - 1).read(queue);
192 } //end compute namespace
193 } //end boost namespace
195 #endif // BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP