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_INCLUDES_HPP
12 #define BOOST_COMPUTE_ALGORITHM_INCLUDES_HPP
13
14 #include <iterator>
15
16 #include <boost/static_assert.hpp>
17
18 #include <boost/compute/algorithm/detail/balanced_path.hpp>
19 #include <boost/compute/algorithm/fill_n.hpp>
20 #include <boost/compute/algorithm/find.hpp>
21 #include <boost/compute/container/vector.hpp>
22 #include <boost/compute/detail/iterator_range_size.hpp>
23 #include <boost/compute/detail/meta_kernel.hpp>
24 #include <boost/compute/detail/read_write_single_value.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 includes kernel class
34 ///
35 /// Subclass of meta_kernel to perform includes operation after tiling
36 ///
37 class serial_includes_kernel : meta_kernel
38 {
39 public:
40
serial_includes_kernel()41 serial_includes_kernel() : meta_kernel("includes")
42 {
43
44 }
45
46 template<class InputIterator1, class InputIterator2,
47 class InputIterator3, class InputIterator4,
48 class OutputIterator>
set_range(InputIterator1 first1,InputIterator2 first2,InputIterator3 tile_first1,InputIterator3 tile_last1,InputIterator4 tile_first2,OutputIterator result)49 void set_range(InputIterator1 first1,
50 InputIterator2 first2,
51 InputIterator3 tile_first1,
52 InputIterator3 tile_last1,
53 InputIterator4 tile_first2,
54 OutputIterator result)
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 includes = 1;\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 " start1++;\n" <<
75 " else\n" <<
76 " {\n" <<
77 " includes = 0;\n" <<
78 " break;\n" <<
79 " }\n" <<
80 "}\n" <<
81 "if(start2<end2)\n" <<
82 " includes = 0;\n" <<
83 result[expr<uint_>("i")] << " = includes;\n";
84 }
85
exec(command_queue & queue)86 event exec(command_queue &queue)
87 {
88 if(m_count == 0) {
89 return event();
90 }
91
92 return exec_1d(queue, 0, m_count);
93 }
94
95 private:
96 size_t m_count;
97 };
98
99 } //end detail namespace
100
101 ///
102 /// \brief Includes algorithm
103 ///
104 /// Finds if the sorted range [first1, last1) includes the sorted
105 /// range [first2, last2). In other words, it checks if [first1, last1) is
106 /// a superset of [first2, last2).
107 ///
108 /// \return True, if [first1, last1) includes [first2, last2). False otherwise.
109 ///
110 /// \param first1 Iterator pointing to start of first set
111 /// \param last1 Iterator pointing to end of first set
112 /// \param first2 Iterator pointing to start of second set
113 /// \param last2 Iterator pointing to end of second set
114 /// \param queue Queue on which to execute
115 ///
116 /// Space complexity: \Omega(distance(\p first1, \p last1) + distance(\p first2, \p last2))
117 template<class InputIterator1, class InputIterator2>
includes(InputIterator1 first1,InputIterator1 last1,InputIterator2 first2,InputIterator2 last2,command_queue & queue=system::default_queue ())118 inline bool includes(InputIterator1 first1,
119 InputIterator1 last1,
120 InputIterator2 first2,
121 InputIterator2 last2,
122 command_queue &queue = system::default_queue())
123 {
124 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
125 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
126
127 size_t tile_size = 1024;
128
129 size_t count1 = detail::iterator_range_size(first1, last1);
130 size_t count2 = detail::iterator_range_size(first2, last2);
131
132 vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
133 vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
134
135 // Tile the sets
136 detail::balanced_path_kernel tiling_kernel;
137 tiling_kernel.tile_size = static_cast<unsigned int>(tile_size);
138 tiling_kernel.set_range(first1, last1, first2, last2,
139 tile_a.begin()+1, tile_b.begin()+1);
140 fill_n(tile_a.begin(), 1, uint_(0), queue);
141 fill_n(tile_b.begin(), 1, uint_(0), queue);
142 tiling_kernel.exec(queue);
143
144 fill_n(tile_a.end()-1, 1, static_cast<uint_>(count1), queue);
145 fill_n(tile_b.end()-1, 1, static_cast<uint_>(count2), queue);
146
147 vector<uint_> result((count1+count2+tile_size-1)/tile_size, queue.get_context());
148
149 // Find individually
150 detail::serial_includes_kernel includes_kernel;
151 includes_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(),
152 tile_b.begin(), result.begin());
153
154 includes_kernel.exec(queue);
155
156 return find(result.begin(), result.end(), 0, queue) == result.end();
157 }
158
159 } //end compute namespace
160 } //end boost namespace
161
162 #endif // BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP
163