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