ceph/src/boost/boost/compute/algorithm/find_end.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_FIND_END_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_FIND_END_HPP
  13
  14 #include <boost/static_assert.hpp>
  15
  16 #include <boost/compute/algorithm/copy.hpp>
  17 #include <boost/compute/algorithm/detail/search_all.hpp>
  18 #include <boost/compute/container/detail/scalar.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/system.hpp>
  23 #include <boost/compute/type_traits/is_device_iterator.hpp>
  24
  25 namespace boost {
  26 namespace compute {
  27 namespace detail {
  28
  29 ///
  30 /// \brief Helper function for find_end
  31 ///
  32 /// Basically a copy of find_if which returns last occurrence
  33 /// instead of first occurrence
  34 ///
  35 template<class InputIterator, class UnaryPredicate>
  36 inline InputIterator find_end_helper(InputIterator first,
  37                                      InputIterator last,
  38                                      UnaryPredicate predicate,
  39                                      command_queue &queue)
  40 {
  41     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
  42     typedef typename std::iterator_traits<InputIterator>::difference_type difference_type;
  43
  44     size_t count = detail::iterator_range_size(first, last);
  45     if(count == 0){
  46         return last;
  47     }
  48
  49     const context &context = queue.get_context();
  50
  51     detail::meta_kernel k("find_end");
  52     size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
  53     atomic_max<int_> atomic_max_int;
  54
  55     k << k.decl<const int_>("i") << " = get_global_id(0);\n"
  56       << k.decl<const value_type>("value") << "="
  57       <<     first[k.var<const int_>("i")] << ";\n"
  58       << "if(" << predicate(k.var<const value_type>("value")) << "){\n"
  59       << "    " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
  60       << "}\n";
  61
  62     kernel kernel = k.compile(context);
  63
  64     scalar<int_> index(context);
  65     kernel.set_arg(index_arg, index.get_buffer());
  66
  67     index.write(static_cast<int_>(-1), queue);
  68
  69     queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
  70
  71     int result = static_cast<int>(index.read(queue));
  72
  73     if(result == -1){
  74         return last;
  75     }
  76     else {
  77         return first + static_cast<difference_type>(result);
  78     }
  79 }
  80
  81 } // end detail namespace
  82
  83 ///
  84 /// \brief Substring matching algorithm
  85 ///
  86 /// Searches for the last match of the pattern [p_first, p_last)
  87 /// in text [t_first, t_last).
  88 /// \return Iterator pointing to beginning of last occurence
  89 ///
  90 /// \param t_first Iterator pointing to start of text
  91 /// \param t_last Iterator pointing to end of text
  92 /// \param p_first Iterator pointing to start of pattern
  93 /// \param p_last Iterator pointing to end of pattern
  94 /// \param queue Queue on which to execute
  95 ///
  96 /// Space complexity: \Omega(n)
  97 ///
  98 template<class TextIterator, class PatternIterator>
  99 inline TextIterator find_end(TextIterator t_first,
 100                              TextIterator t_last,
 101                              PatternIterator p_first,
 102                              PatternIterator p_last,
 103                              command_queue &queue = system::default_queue())
 104 {
 105     BOOST_STATIC_ASSERT(is_device_iterator<TextIterator>::value);
 106     BOOST_STATIC_ASSERT(is_device_iterator<PatternIterator>::value);
 107
 108     const context &context = queue.get_context();
 109
 110     // there is no need to check if pattern starts at last n - 1 indices
 111     vector<uint_> matching_indices(
 112         detail::iterator_range_size(t_first, t_last)
 113             + 1 - detail::iterator_range_size(p_first, p_last),
 114         context
 115     );
 116
 117     detail::search_kernel<PatternIterator,
 118                           TextIterator,
 119                           vector<uint_>::iterator> kernel;
 120
 121     kernel.set_range(p_first, p_last, t_first, t_last, matching_indices.begin());
 122     kernel.exec(queue);
 123
 124     using boost::compute::_1;
 125
 126     vector<uint_>::iterator index =
 127         detail::find_end_helper(
 128             matching_indices.begin(),
 129             matching_indices.end(),
 130             _1 == 1,
 131             queue
 132         );
 133
 134     // pattern was not found
 135     if(index == matching_indices.end())
 136         return t_last;
 137
 138     return t_first + detail::iterator_range_size(matching_indices.begin(), index);
 139 }
 140
 141 } //end compute namespace
 142 } //end boost namespace
 143
 144 #endif // BOOST_COMPUTE_ALGORITHM_FIND_END_HPP