123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179 |
- #ifndef BOOST_COMPUTE_ALGORITHM_SET_INTERSECTION_HPP
- #define BOOST_COMPUTE_ALGORITHM_SET_INTERSECTION_HPP
- #include <iterator>
- #include <boost/static_assert.hpp>
- #include <boost/compute/algorithm/detail/compact.hpp>
- #include <boost/compute/algorithm/detail/balanced_path.hpp>
- #include <boost/compute/algorithm/exclusive_scan.hpp>
- #include <boost/compute/algorithm/fill_n.hpp>
- #include <boost/compute/container/vector.hpp>
- #include <boost/compute/detail/iterator_range_size.hpp>
- #include <boost/compute/detail/meta_kernel.hpp>
- #include <boost/compute/system.hpp>
- #include <boost/compute/type_traits/is_device_iterator.hpp>
- namespace boost {
- namespace compute {
- namespace detail {
- class serial_set_intersection_kernel : meta_kernel
- {
- public:
- unsigned int tile_size;
- serial_set_intersection_kernel() : meta_kernel("set_intersection")
- {
- tile_size = 4;
- }
- template<class InputIterator1, class InputIterator2,
- class InputIterator3, class InputIterator4,
- class OutputIterator1, class OutputIterator2>
- void set_range(InputIterator1 first1,
- InputIterator2 first2,
- InputIterator3 tile_first1,
- InputIterator3 tile_last1,
- InputIterator4 tile_first2,
- OutputIterator1 result,
- OutputIterator2 counts)
- {
- m_count = iterator_range_size(tile_first1, tile_last1) - 1;
- *this <<
- "uint i = get_global_id(0);\n" <<
- "uint start1 = " << tile_first1[expr<uint_>("i")] << ";\n" <<
- "uint end1 = " << tile_first1[expr<uint_>("i+1")] << ";\n" <<
- "uint start2 = " << tile_first2[expr<uint_>("i")] << ";\n" <<
- "uint end2 = " << tile_first2[expr<uint_>("i+1")] << ";\n" <<
- "uint index = i*" << tile_size << ";\n" <<
- "uint count = 0;\n" <<
- "while(start1<end1 && start2<end2)\n" <<
- "{\n" <<
- " if(" << first1[expr<uint_>("start1")] << " == " <<
- first2[expr<uint_>("start2")] << ")\n" <<
- " {\n" <<
- result[expr<uint_>("index")] <<
- " = " << first1[expr<uint_>("start1")] << ";\n" <<
- " index++; count++;\n" <<
- " start1++; start2++;\n" <<
- " }\n" <<
- " else if(" << first1[expr<uint_>("start1")] << " < " <<
- first2[expr<uint_>("start2")] << ")\n" <<
- " start1++;\n" <<
- " else start2++;\n" <<
- "}\n" <<
- counts[expr<uint_>("i")] << " = count;\n";
- }
- event exec(command_queue &queue)
- {
- if(m_count == 0) {
- return event();
- }
- return exec_1d(queue, 0, m_count);
- }
- private:
- size_t m_count;
- };
- }
- template<class InputIterator1, class InputIterator2, class OutputIterator>
- inline OutputIterator set_intersection(InputIterator1 first1,
- InputIterator1 last1,
- InputIterator2 first2,
- InputIterator2 last2,
- OutputIterator result,
- command_queue &queue = system::default_queue())
- {
- BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
- BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
- BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
- typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
- int tile_size = 1024;
- int count1 = detail::iterator_range_size(first1, last1);
- int count2 = detail::iterator_range_size(first2, last2);
- vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
- vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
-
- detail::balanced_path_kernel tiling_kernel;
- tiling_kernel.tile_size = tile_size;
- tiling_kernel.set_range(first1, last1, first2, last2,
- tile_a.begin()+1, tile_b.begin()+1);
- fill_n(tile_a.begin(), 1, 0, queue);
- fill_n(tile_b.begin(), 1, 0, queue);
- tiling_kernel.exec(queue);
- fill_n(tile_a.end()-1, 1, count1, queue);
- fill_n(tile_b.end()-1, 1, count2, queue);
- vector<value_type> temp_result(count1+count2, queue.get_context());
- vector<uint_> counts((count1+count2+tile_size-1)/tile_size + 1, queue.get_context());
- fill_n(counts.end()-1, 1, 0, queue);
-
- detail::serial_set_intersection_kernel intersection_kernel;
- intersection_kernel.tile_size = tile_size;
- intersection_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(),
- tile_b.begin(), temp_result.begin(), counts.begin());
- intersection_kernel.exec(queue);
- exclusive_scan(counts.begin(), counts.end(), counts.begin(), queue);
-
- detail::compact_kernel compact_kernel;
- compact_kernel.tile_size = tile_size;
- compact_kernel.set_range(temp_result.begin(), counts.begin(), counts.end(), result);
- compact_kernel.exec(queue);
- return result + (counts.end() - 1).read(queue);
- }
- }
- }
- #endif
|