ceph/src/boost/boost/compute/algorithm/set_union.hpp

   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_UNION_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_SET_UNION_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/detail/compact.hpp>
  20 #include <boost/compute/algorithm/exclusive_scan.hpp>
  21 #include <boost/compute/algorithm/fill_n.hpp>
  22 #include <boost/compute/container/vector.hpp>
  23 #include <boost/compute/detail/iterator_range_size.hpp>
  24 #include <boost/compute/detail/meta_kernel.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 set union kernel class
  34 ///
  35 /// Subclass of meta_kernel to perform serial set union after tiling
  36 ///
  37 class serial_set_union_kernel : meta_kernel
  38 {
  39 public:
  40     unsigned int tile_size;
  41
  42     serial_set_union_kernel() : meta_kernel("set_union")
  43     {
  44         tile_size = 4;
  45     }
  46
  47     template<class InputIterator1, class InputIterator2,
  48              class InputIterator3, class InputIterator4,
  49              class OutputIterator1, class OutputIterator2>
  50     void set_range(InputIterator1 first1,
  51                     InputIterator2 first2,
  52                     InputIterator3 tile_first1,
  53                     InputIterator3 tile_last1,
  54                     InputIterator4 tile_first2,
  55                     OutputIterator1 result,
  56                     OutputIterator2 counts)
  57     {
  58         m_count = iterator_range_size(tile_first1, tile_last1) - 1;
  59
  60         *this <<
  61         "uint i = get_global_id(0);\n" <<
  62         "uint start1 = " << tile_first1[expr<uint_>("i")] << ";\n" <<
  63         "uint end1 = " << tile_first1[expr<uint_>("i+1")] << ";\n" <<
  64         "uint start2 = " << tile_first2[expr<uint_>("i")] << ";\n" <<
  65         "uint end2 = " << tile_first2[expr<uint_>("i+1")] << ";\n" <<
  66         "uint index = i*" << tile_size << ";\n" <<
  67         "uint count = 0;\n" <<
  68         "while(start1<end1 && start2<end2)\n" <<
  69         "{\n" <<
  70         "   if(" << first1[expr<uint_>("start1")] << " == " <<
  71                     first2[expr<uint_>("start2")] << ")\n" <<
  72         "   {\n" <<
  73                 result[expr<uint_>("index")] <<
  74                     " = " << first1[expr<uint_>("start1")] << ";\n" <<
  75         "       index++; count++;\n" <<
  76         "       start1++; start2++;\n" <<
  77         "   }\n" <<
  78         "   else if(" << first1[expr<uint_>("start1")] << " < " <<
  79                         first2[expr<uint_>("start2")] << ")\n" <<
  80         "   {\n" <<
  81                 result[expr<uint_>("index")] <<
  82                     " = " << first1[expr<uint_>("start1")] << ";\n" <<
  83         "       index++; count++;\n" <<
  84         "       start1++;\n" <<
  85         "   }\n" <<
  86         "   else\n" <<
  87         "   {\n" <<
  88                 result[expr<uint_>("index")] <<
  89                     " = " << first2[expr<uint_>("start2")] << ";\n" <<
  90         "       index++; count++;\n" <<
  91         "       start2++;\n" <<
  92         "   }\n" <<
  93         "}\n" <<
  94         "while(start1<end1)\n" <<
  95         "{\n" <<
  96             result[expr<uint_>("index")] <<
  97                 " = " << first1[expr<uint_>("start1")] << ";\n" <<
  98         "   index++; count++;\n" <<
  99         "   start1++;\n" <<
 100         "}\n" <<
 101         "while(start2<end2)\n" <<
 102         "{\n" <<
 103             result[expr<uint_>("index")] <<
 104                 " = " << first2[expr<uint_>("start2")] << ";\n" <<
 105         "   index++; count++;\n" <<
 106         "   start2++;\n" <<
 107         "}\n" <<
 108         counts[expr<uint_>("i")] << " = count;\n";
 109     }
 110
 111     event exec(command_queue &queue)
 112     {
 113         if(m_count == 0) {
 114             return event();
 115         }
 116
 117         return exec_1d(queue, 0, m_count);
 118     }
 119
 120 private:
 121     size_t m_count;
 122 };
 123
 124 } //end detail namespace
 125
 126 ///
 127 /// \brief Set union algorithm
 128 ///
 129 /// Finds the union of the sorted range [first1, last1) with the sorted
 130 /// range [first2, last2) and stores it in range starting at result
 131 /// \return Iterator pointing to end of union
 132 ///
 133 /// \param first1 Iterator pointing to start of first set
 134 /// \param last1 Iterator pointing to end of first set
 135 /// \param first2 Iterator pointing to start of second set
 136 /// \param last2 Iterator pointing to end of second set
 137 /// \param result Iterator pointing to start of range in which the union
 138 /// will be stored
 139 /// \param queue Queue on which to execute
 140 ///
 141 /// Space complexity:
 142 /// \Omega(2(distance(\p first1, \p last1) + distance(\p first2, \p last2)))
 143 template<class InputIterator1, class InputIterator2, class OutputIterator>
 144 inline OutputIterator set_union(InputIterator1 first1,
 145                                 InputIterator1 last1,
 146                                 InputIterator2 first2,
 147                                 InputIterator2 last2,
 148                                 OutputIterator result,
 149                                 command_queue &queue = system::default_queue())
 150 {
 151     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
 152     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
 153     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
 154
 155     typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
 156
 157     int tile_size = 1024;
 158
 159     int count1 = detail::iterator_range_size(first1, last1);
 160     int count2 = detail::iterator_range_size(first2, last2);
 161
 162     vector<uint_> tile_a((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
 163     vector<uint_> tile_b((count1+count2+tile_size-1)/tile_size+1, queue.get_context());
 164
 165     // Tile the sets
 166     detail::balanced_path_kernel tiling_kernel;
 167     tiling_kernel.tile_size = tile_size;
 168     tiling_kernel.set_range(first1, last1, first2, last2,
 169                             tile_a.begin()+1, tile_b.begin()+1);
 170     fill_n(tile_a.begin(), 1, 0, queue);
 171     fill_n(tile_b.begin(), 1, 0, queue);
 172     tiling_kernel.exec(queue);
 173
 174     fill_n(tile_a.end()-1, 1, count1, queue);
 175     fill_n(tile_b.end()-1, 1, count2, queue);
 176
 177     vector<value_type> temp_result(count1+count2, queue.get_context());
 178     vector<uint_> counts((count1+count2+tile_size-1)/tile_size + 1, queue.get_context());
 179     fill_n(counts.end()-1, 1, 0, queue);
 180
 181     // Find individual unions
 182     detail::serial_set_union_kernel union_kernel;
 183     union_kernel.tile_size = tile_size;
 184     union_kernel.set_range(first1, first2, tile_a.begin(), tile_a.end(),
 185                                   tile_b.begin(), temp_result.begin(), counts.begin());
 186
 187     union_kernel.exec(queue);
 188
 189     exclusive_scan(counts.begin(), counts.end(), counts.begin(), queue);
 190
 191     // Compact the results
 192     detail::compact_kernel compact_kernel;
 193     compact_kernel.tile_size = tile_size;
 194     compact_kernel.set_range(temp_result.begin(), counts.begin(), counts.end(), result);
 195
 196     compact_kernel.exec(queue);
 197
 198     return result + (counts.end() - 1).read(queue);
 199 }
 200
 201 } //end compute namespace
 202 } //end boost namespace
 203
 204 #endif // BOOST_COMPUTE_ALGORITHM_SET_UNION_HPP