ceph/src/boost/boost/compute/algorithm/detail/merge_path.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_DETAIL_MERGE_PATH_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_DETAIL_MERGE_PATH_HPP
  13
  14 #include <iterator>
  15
  16 #include <boost/compute/algorithm/find_if.hpp>
  17 #include <boost/compute/container/vector.hpp>
  18 #include <boost/compute/detail/iterator_range_size.hpp>
  19 #include <boost/compute/detail/meta_kernel.hpp>
  20 #include <boost/compute/lambda.hpp>
  21 #include <boost/compute/system.hpp>
  22
  23 namespace boost {
  24 namespace compute {
  25 namespace detail {
  26
  27 ///
  28 /// \brief Merge Path kernel class
  29 ///
  30 /// Subclass of meta_kernel to break two sets into tiles according
  31 /// to their merge path
  32 ///
  33 class merge_path_kernel : public meta_kernel
  34 {
  35 public:
  36     unsigned int tile_size;
  37
  38     merge_path_kernel() : meta_kernel("merge_path")
  39     {
  40         tile_size = 4;
  41     }
  42
  43     template<class InputIterator1, class InputIterator2,
  44              class OutputIterator1, class OutputIterator2,
  45              class Compare>
  46     void set_range(InputIterator1 first1,
  47                    InputIterator1 last1,
  48                    InputIterator2 first2,
  49                    InputIterator2 last2,
  50                    OutputIterator1 result_a,
  51                    OutputIterator2 result_b,
  52                    Compare comp)
  53     {
  54         m_a_count = iterator_range_size(first1, last1);
  55         m_a_count_arg = add_arg<uint_>("a_count");
  56
  57         m_b_count = iterator_range_size(first2, last2);
  58         m_b_count_arg = add_arg<uint_>("b_count");
  59
  60         *this <<
  61             "uint i = get_global_id(0);\n" <<
  62             "uint target = (i+1)*" << tile_size << ";\n" <<
  63             "uint start = max(convert_int(0),convert_int(target)-convert_int(b_count));\n" <<
  64             "uint end = min(target,a_count);\n" <<
  65             "uint a_index, b_index;\n" <<
  66             "while(start<end)\n" <<
  67             "{\n" <<
  68             "   a_index = (start + end)/2;\n" <<
  69             "   b_index = target - a_index - 1;\n" <<
  70             "   if(!(" << comp(first2[expr<uint_>("b_index")],
  71                               first1[expr<uint_>("a_index")]) << "))\n" <<
  72             "       start = a_index + 1;\n" <<
  73             "   else end = a_index;\n" <<
  74             "}\n" <<
  75             result_a[expr<uint_>("i")] << " = start;\n" <<
  76             result_b[expr<uint_>("i")] << " = target - start;\n";
  77     }
  78
  79     template<class InputIterator1, class InputIterator2,
  80              class OutputIterator1, class OutputIterator2>
  81     void set_range(InputIterator1 first1,
  82                    InputIterator1 last1,
  83                    InputIterator2 first2,
  84                    InputIterator2 last2,
  85                    OutputIterator1 result_a,
  86                    OutputIterator2 result_b)
  87     {
  88         typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
  89         ::boost::compute::less<value_type> less_than;
  90         set_range(first1, last1, first2, last2, result_a, result_b, less_than);
  91     }
  92
  93     event exec(command_queue &queue)
  94     {
  95         if((m_a_count + m_b_count)/tile_size == 0) {
  96             return event();
  97         }
  98
  99         set_arg(m_a_count_arg, uint_(m_a_count));
 100         set_arg(m_b_count_arg, uint_(m_b_count));
 101
 102         return exec_1d(queue, 0, (m_a_count + m_b_count)/tile_size);
 103     }
 104
 105 private:
 106     size_t m_a_count;
 107     size_t m_a_count_arg;
 108     size_t m_b_count;
 109     size_t m_b_count_arg;
 110 };
 111
 112 } //end detail namespace
 113 } //end compute namespace
 114 } //end boost namespace
 115
 116 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_MERGE_PATH_HPP