ceph/src/boost/boost/compute/algorithm/prev_permutation.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_PREV_PERMUTATION_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_PREV_PERMUTATION_HPP
  13
  14 #include <iterator>
  15
  16 #include <boost/compute/system.hpp>
  17 #include <boost/compute/command_queue.hpp>
  18 #include <boost/compute/container/detail/scalar.hpp>
  19 #include <boost/compute/algorithm/reverse.hpp>
  20
  21 namespace boost {
  22 namespace compute {
  23 namespace detail {
  24
  25 ///
  26 /// \brief Helper function for prev_permutation
  27 ///
  28 /// To find rightmost element which is greater
  29 /// than its next element
  30 ///
  31 template<class InputIterator>
  32 inline InputIterator prev_permutation_helper(InputIterator first,
  33                                              InputIterator last,
  34                                              command_queue &queue)
  35 {
  36     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
  37
  38     size_t count = detail::iterator_range_size(first, last);
  39     if(count == 0 || count == 1){
  40         return last;
  41     }
  42     count = count - 1;
  43     const context &context = queue.get_context();
  44
  45     detail::meta_kernel k("prev_permutation");
  46     size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
  47     atomic_max<int_> atomic_max_int;
  48
  49     k << k.decl<const int_>("i") << " = get_global_id(0);\n"
  50       << k.decl<const value_type>("cur_value") << "="
  51       <<     first[k.var<const int_>("i")] << ";\n"
  52       << k.decl<const value_type>("next_value") << "="
  53       <<     first[k.expr<const int_>("i+1")] << ";\n"
  54       << "if(cur_value > next_value){\n"
  55       << "    " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
  56       << "}\n";
  57
  58     kernel kernel = k.compile(context);
  59
  60     scalar<int_> index(context);
  61     kernel.set_arg(index_arg, index.get_buffer());
  62
  63     index.write(static_cast<int_>(-1), queue);
  64
  65     queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
  66
  67     int result = static_cast<int>(index.read(queue));
  68     if(result == -1) return last;
  69     else return first + result;
  70 }
  71
  72 ///
  73 /// \brief Helper function for prev_permutation
  74 ///
  75 /// To find the largest element to the right of the element found above
  76 /// that is smaller than it
  77 ///
  78 template<class InputIterator, class ValueType>
  79 inline InputIterator pp_floor(InputIterator first,
  80                               InputIterator last,
  81                               ValueType value,
  82                               command_queue &queue)
  83 {
  84     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
  85
  86     size_t count = detail::iterator_range_size(first, last);
  87     if(count == 0){
  88         return last;
  89     }
  90     const context &context = queue.get_context();
  91
  92     detail::meta_kernel k("pp_floor");
  93     size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
  94     size_t value_arg = k.add_arg<value_type>(memory_object::private_memory, "value");
  95     atomic_max<int_> atomic_max_int;
  96
  97     k << k.decl<const int_>("i") << " = get_global_id(0);\n"
  98       << k.decl<const value_type>("cur_value") << "="
  99       <<     first[k.var<const int_>("i")] << ";\n"
 100       << "if(cur_value >= " << first[k.expr<int_>("*index")]
 101       << "      && cur_value < value){\n"
 102       << "    " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
 103       << "}\n";
 104
 105     kernel kernel = k.compile(context);
 106
 107     scalar<int_> index(context);
 108     kernel.set_arg(index_arg, index.get_buffer());
 109
 110     index.write(static_cast<int_>(0), queue);
 111
 112     kernel.set_arg(value_arg, value);
 113
 114     queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
 115
 116     int result = static_cast<int>(index.read(queue));
 117     return first + result;
 118 }
 119
 120 } // end detail namespace
 121
 122 ///
 123 /// \brief Permutation generating algorithm
 124 ///
 125 /// Transforms the range [first, last) into the previous permutation from
 126 /// the set of all permutations arranged in lexicographic order
 127 /// \return Boolean value signifying if the first permutation was crossed
 128 /// and the range was reset
 129 ///
 130 /// \param first Iterator pointing to start of range
 131 /// \param last Iterator pointing to end of range
 132 /// \param queue Queue on which to execute
 133 ///
 134 /// Space complexity: \Omega(1)
 135 template<class InputIterator>
 136 inline bool prev_permutation(InputIterator first,
 137                              InputIterator last,
 138                              command_queue &queue = system::default_queue())
 139 {
 140     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
 141
 142     if(first == last) return false;
 143
 144     InputIterator first_element =
 145         detail::prev_permutation_helper(first, last, queue);
 146
 147     if(first_element == last)
 148     {
 149         reverse(first, last, queue);
 150         return false;
 151     }
 152
 153     value_type first_value = first_element.read(queue);
 154
 155     InputIterator ceiling_element =
 156         detail::pp_floor(first_element + 1, last, first_value, queue);
 157
 158     value_type ceiling_value = ceiling_element.read(queue);
 159
 160     first_element.write(ceiling_value, queue);
 161     ceiling_element.write(first_value, queue);
 162
 163     reverse(first_element + 1, last, queue);
 164
 165     return true;
 166 }
 167
 168 } // end compute namespace
 169 } // end boost namespace
 170
 171 #endif // BOOST_COMPUTE_ALGORITHM_PREV_PERMUTATION_HPP