]>
Commit | Line | Data |
---|---|---|
7c673cae FG |
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_DIFFERENCE_HPP | |
12 | #define BOOST_COMPUTE_ALGORITHM_SET_DIFFERENCE_HPP | |
13 | ||
14 | #include <iterator> | |
15 | ||
16 | #include <boost/compute/algorithm/detail/compact.hpp> | |
17 | #include <boost/compute/algorithm/detail/balanced_path.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> | |
24 | ||
25 | namespace boost { | |
26 | namespace compute { | |
27 | namespace detail { | |
28 | ||
29 | /// | |
30 | /// \brief Serial set difference kernel class | |
31 | /// | |
32 | /// Subclass of meta_kernel to perform serial set difference after tiling | |
33 | /// | |
34 | class serial_set_difference_kernel : meta_kernel | |
35 | { | |
36 | public: | |
37 | unsigned int tile_size; | |
38 | ||
39 | serial_set_difference_kernel() : meta_kernel("set_difference") | |
40 | { | |
41 | tile_size = 4; | |
42 | } | |
43 | ||
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) | |
54 | { | |
55 | m_count = iterator_range_size(tile_first1, tile_last1) - 1; | |
56 | ||
57 | *this << | |
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" << | |
66 | "{\n" << | |
67 | " if(" << first1[expr<uint_>("start1")] << " == " << | |
68 | first2[expr<uint_>("start2")] << ")\n" << | |
69 | " {\n" << | |
70 | " start1++; start2++;\n" << | |
71 | " }\n" << | |
72 | " else if(" << first1[expr<uint_>("start1")] << " < " << | |
73 | first2[expr<uint_>("start2")] << ")\n" << | |
74 | " {\n" << | |
75 | result[expr<uint_>("index")] << | |
76 | " = " << first1[expr<uint_>("start1")] << ";\n" << | |
77 | " index++; count++;\n" << | |
78 | " start1++;\n" << | |
79 | " }\n" << | |
80 | " else\n" << | |
81 | " {\n" << | |
82 | " start2++;\n" << | |
83 | " }\n" << | |
84 | "}\n" << | |
85 | "while(start1<end1)\n" << | |
86 | "{\n" << | |
87 | result[expr<uint_>("index")] << | |
88 | " = " << first1[expr<uint_>("start1")] << ";\n" << | |
89 | " index++; count++;\n" << | |
90 | " start1++;\n" << | |
91 | "}\n" << | |
92 | counts[expr<uint_>("i")] << " = count;\n"; | |
93 | } | |
94 | ||
95 | event exec(command_queue &queue) | |
96 | { | |
97 | if(m_count == 0) { | |
98 | return event(); | |
99 | } | |
100 | ||
101 | return exec_1d(queue, 0, m_count); | |
102 | } | |
103 | ||
104 | private: | |
105 | size_t m_count; | |
106 | }; | |
107 | ||
108 | } //end detail namespace | |
109 | ||
110 | /// | |
111 | /// \brief Set difference algorithm | |
112 | /// | |
113 | /// Finds the difference of the sorted range [first2, last2) from the sorted | |
114 | /// range [first1, last1) and stores it in range starting at result | |
115 | /// \return Iterator pointing to end of difference | |
116 | /// | |
117 | /// \param first1 Iterator pointing to start of first set | |
118 | /// \param last1 Iterator pointing to end of first set | |
119 | /// \param first2 Iterator pointing to start of second set | |
120 | /// \param last2 Iterator pointing to end of second set | |
121 | /// \param result Iterator pointing to start of range in which the difference | |
122 | /// will be stored | |
123 | /// \param queue Queue on which to execute | |
124 | /// | |
125 | template<class InputIterator1, class InputIterator2, class OutputIterator> | |
126 | inline OutputIterator set_difference(InputIterator1 first1, | |
127 | InputIterator1 last1, | |
128 | InputIterator2 first2, | |
129 | InputIterator2 last2, | |
130 | OutputIterator result, | |
131 | command_queue &queue = system::default_queue()) | |
132 | { | |
133 | typedef typename std::iterator_traits<InputIterator1>::value_type value_type; | |
134 | ||
135 | int tile_size = 1024; | |
136 | ||
137 | int count1 = detail::iterator_range_size(first1, last1); | |
138 | int count2 = detail::iterator_range_size(first2, last2); | |
139 | ||
140 | vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context()); | |
141 | vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context()); | |
142 | ||
143 | // Tile the sets | |
144 | detail::balanced_path_kernel tiling_kernel; | |
145 | tiling_kernel.tile_size = tile_size; | |
146 | tiling_kernel.set_range(first1, last1, first2, last2, | |
147 | tile_a.begin()+1, tile_b.begin()+1); | |
148 | fill_n(tile_a.begin(), 1, 0, queue); | |
149 | fill_n(tile_b.begin(), 1, 0, queue); | |
150 | tiling_kernel.exec(queue); | |
151 | ||
152 | fill_n(tile_a.end()-1, 1, count1, queue); | |
153 | fill_n(tile_b.end()-1, 1, count2, queue); | |
154 | ||
155 | vector<value_type> temp_result(count1+count2, queue.get_context()); | |
156 | vector<uint_> counts((count1+count2+tile_size-1)/tile_size + 1, queue.get_context()); | |
157 | fill_n(counts.end()-1, 1, 0, queue); | |
158 | ||
159 | // Find individual differences | |
160 | detail::serial_set_difference_kernel difference_kernel; | |
161 | difference_kernel.tile_size = tile_size; | |
162 | difference_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(), | |
163 | tile_b.begin(), temp_result.begin(), counts.begin()); | |
164 | ||
165 | difference_kernel.exec(queue); | |
166 | ||
167 | exclusive_scan(counts.begin(), counts.end(), counts.begin(), queue); | |
168 | ||
169 | // Compact the results | |
170 | detail::compact_kernel compact_kernel; | |
171 | compact_kernel.tile_size = tile_size; | |
172 | compact_kernel.set_range(temp_result.begin(), counts.begin(), counts.end(), result); | |
173 | ||
174 | compact_kernel.exec(queue); | |
175 | ||
176 | return result + (counts.end() - 1).read(queue); | |
177 | } | |
178 | ||
179 | } //end compute namespace | |
180 | } //end boost namespace | |
181 | ||
182 | #endif // BOOST_COMPUTE_ALGORITHM_SET_DIFFERENCE_HPP |