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