compute/algorithm/set_union.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2014 Roshan <thisisroshansmail@gmail.com>
0003 //
0004 // Distributed under the Boost Software License, Version 1.0
0005 // See accompanying file LICENSE_1_0.txt or copy at
0006 // http://www.boost.org/LICENSE_1_0.txt
0007 //
0008 // See http://boostorg.github.com/compute for more information.
0009 //---------------------------------------------------------------------------//
0010
0011 #ifndef BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP
0013
0014 #include <iterator>
0015
0016 #include <boost/static_assert.hpp>
0017
0018 #include <boost/compute/algorithm/detail/balanced_path.hpp>
0019 #include <boost/compute/algorithm/detail/compact.hpp>
0020 #include <boost/compute/algorithm/exclusive_scan.hpp>
0021 #include <boost/compute/algorithm/fill_n.hpp>
0022 #include <boost/compute/container/vector.hpp>
0023 #include <boost/compute/detail/iterator_range_size.hpp>
0024 #include <boost/compute/detail/meta_kernel.hpp>
0025 #include <boost/compute/system.hpp>
0026 #include <boost/compute/type_traits/is_device_iterator.hpp>
0027
0028 namespace boost {
0029 namespace compute {
0030 namespace detail {
0031
0032 ///
0033 /// \brief Serial set union kernel class
0034 ///
0035 /// Subclass of meta_kernel to perform serial set union after tiling
0036 ///
0037 class serial_set_union_kernel : meta_kernel
0038 {
0039 public:
0040     unsigned int tile_size;
0041
0042     serial_set_union_kernel() : meta_kernel("set_union")
0043     {
0044         tile_size = 4;
0045     }
0046
0047     template<class InputIterator1, class InputIterator2,
0048              class InputIterator3, class InputIterator4,
0049              class OutputIterator1, class OutputIterator2>
0050     void set_range(InputIterator1 first1,
0051                     InputIterator2 first2,
0052                     InputIterator3 tile_first1,
0053                     InputIterator3 tile_last1,
0054                     InputIterator4 tile_first2,
0055                     OutputIterator1 result,
0056                     OutputIterator2 counts)
0057     {
0058         m_count = iterator_range_size(tile_first1, tile_last1) - 1;
0059
0060         *this <<
0061         "uint i = get_global_id(0);\n" <<
0062         "uint start1 = " << tile_first1[expr<uint_>("i")] << ";\n" <<
0063         "uint end1 = " << tile_first1[expr<uint_>("i+1")] << ";\n" <<
0064         "uint start2 = " << tile_first2[expr<uint_>("i")] << ";\n" <<
0065         "uint end2 = " << tile_first2[expr<uint_>("i+1")] << ";\n" <<
0066         "uint index = i*" << tile_size << ";\n" <<
0067         "uint count = 0;\n" <<
0068         "while(start1<end1 && start2<end2)\n" <<
0069         "{\n" <<
0070         "   if(" << first1[expr<uint_>("start1")] << " == " <<
0071                     first2[expr<uint_>("start2")] << ")\n" <<
0072         "   {\n" <<
0073                 result[expr<uint_>("index")] <<
0074                     " = " << first1[expr<uint_>("start1")] << ";\n" <<
0075         "       index++; count++;\n" <<
0076         "       start1++; start2++;\n" <<
0077         "   }\n" <<
0078         "   else if(" << first1[expr<uint_>("start1")] << " < " <<
0079                         first2[expr<uint_>("start2")] << ")\n" <<
0080         "   {\n" <<
0081                 result[expr<uint_>("index")] <<
0082                     " = " << first1[expr<uint_>("start1")] << ";\n" <<
0083         "       index++; count++;\n" <<
0084         "       start1++;\n" <<
0085         "   }\n" <<
0086         "   else\n" <<
0087         "   {\n" <<
0088                 result[expr<uint_>("index")] <<
0089                     " = " << first2[expr<uint_>("start2")] << ";\n" <<
0090         "       index++; count++;\n" <<
0091         "       start2++;\n" <<
0092         "   }\n" <<
0093         "}\n" <<
0094         "while(start1<end1)\n" <<
0095         "{\n" <<
0096             result[expr<uint_>("index")] <<
0097                 " = " << first1[expr<uint_>("start1")] << ";\n" <<
0098         "   index++; count++;\n" <<
0099         "   start1++;\n" <<
0100         "}\n" <<
0101         "while(start2<end2)\n" <<
0102         "{\n" <<
0103             result[expr<uint_>("index")] <<
0104                 " = " << first2[expr<uint_>("start2")] << ";\n" <<
0105         "   index++; count++;\n" <<
0106         "   start2++;\n" <<
0107         "}\n" <<
0108         counts[expr<uint_>("i")] << " = count;\n";
0109     }
0110
0111     event exec(command_queue &queue)
0112     {
0113         if(m_count == 0) {
0114             return event();
0115         }
0116
0117         return exec_1d(queue, 0, m_count);
0118     }
0119
0120 private:
0121     size_t m_count;
0122 };
0123
0124 } //end detail namespace
0125
0126 ///
0127 /// \brief Set union algorithm
0128 ///
0129 /// Finds the union of the sorted range [first1, last1) with the sorted
0130 /// range [first2, last2) and stores it in range starting at result
0131 /// \return Iterator pointing to end of union
0132 ///
0133 /// \param first1 Iterator pointing to start of first set
0134 /// \param last1 Iterator pointing to end of first set
0135 /// \param first2 Iterator pointing to start of second set
0136 /// \param last2 Iterator pointing to end of second set
0137 /// \param result Iterator pointing to start of range in which the union
0138 /// will be stored
0139 /// \param queue Queue on which to execute
0140 ///
0141 /// Space complexity:
0142 /// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
0143 template<class InputIterator1, class InputIterator2, class OutputIterator>
0144 inline OutputIterator set_union(InputIterator1 first1,
0145                                 InputIterator1 last1,
0146                                 InputIterator2 first2,
0147                                 InputIterator2 last2,
0148                                 OutputIterator result,
0149                                 command_queue &queue = system::default_queue())
0150 {
0151     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
0152     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
0153     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
0154
0155     typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
0156
0157     int tile_size = 1024;
0158
0159     int count1 = detail::iterator_range_size(first1, last1);
0160     int count2 = detail::iterator_range_size(first2, last2);
0161
0162     vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
0163     vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
0164
0165     // Tile the sets
0166     detail::balanced_path_kernel tiling_kernel;
0167     tiling_kernel.tile_size = tile_size;
0168     tiling_kernel.set_range(first1, last1, first2, last2,
0169                             tile_a.begin()+1, tile_b.begin()+1);
0170     fill_n(tile_a.begin(), 1, 0, queue);
0171     fill_n(tile_b.begin(), 1, 0, queue);
0172     tiling_kernel.exec(queue);
0173
0174     fill_n(tile_a.end()-1, 1, count1, queue);
0175     fill_n(tile_b.end()-1, 1, count2, queue);
0176
0177     vector<value_type> temp_result(count1+count2, queue.get_context());
0178     vector<uint_> counts((count1+count2+tile_size-1)/tile_size + 1, queue.get_context());
0179     fill_n(counts.end()-1, 1, 0, queue);
0180
0181     // Find individual unions
0182     detail::serial_set_union_kernel union_kernel;
0183     union_kernel.tile_size = tile_size;
0184     union_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(),
0185                                   tile_b.begin(), temp_result.begin(), counts.begin());
0186
0187     union_kernel.exec(queue);
0188
0189     exclusive_scan(counts.begin(), counts.end(), counts.begin(), queue);
0190
0191     // Compact the results
0192     detail::compact_kernel compact_kernel;
0193     compact_kernel.tile_size = tile_size;
0194     compact_kernel.set_range(temp_result.begin(), counts.begin(), counts.end(), result);
0195
0196     compact_kernel.exec(queue);
0197
0198     return result + (counts.end() - 1).read(queue);
0199 }
0200
0201 } //end compute namespace
0202 } //end boost namespace
0203
0204 #endif // BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP