ceph/src/boost/libs/compute/include/boost/compute/algorithm/next_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_NEXT_PERMUTATION_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_NEXT_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 next_permutation
  27 ///
  28 /// To find rightmost element which is smaller
  29 /// than its next element
  30 ///
  31 template<class InputIterator>
  32 inline InputIterator next_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("next_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 next_permutation
  74 ///
  75 /// To find the smallest element to the right of the element found above
  76 /// that is greater than it
  77 ///
  78 template<class InputIterator, class ValueType>
  79 inline InputIterator np_ceiling(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("np_ceiling");
  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 next permutation from the
 126 /// set of all permutations arranged in lexicographic order
 127 /// \return Boolean value signifying if the last 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 template<class InputIterator>
 135 inline bool next_permutation(InputIterator first,
 136                              InputIterator last,
 137                              command_queue &queue = system::default_queue())
 138 {
 139     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
 140
 141     if(first == last) return false;
 142
 143     InputIterator first_element =
 144         detail::next_permutation_helper(first, last, queue);
 145
 146     if(first_element == last)
 147     {
 148         reverse(first, last, queue);
 149         return false;
 150     }
 151
 152     value_type first_value = first_element.read(queue);
 153
 154     InputIterator ceiling_element =
 155         detail::np_ceiling(first_element + 1, last, first_value, queue);
 156
 157     value_type ceiling_value = ceiling_element.read(queue);
 158
 159     first_element.write(ceiling_value, queue);
 160     ceiling_element.write(first_value, queue);
 161
 162     reverse(first_element + 1, last, queue);
 163
 164     return true;
 165 }
 166
 167 } // end compute namespace
 168 } // end boost namespace
 169
 170 #endif // BOOST_COMPUTE_ALGORITHM_NEXT_PERMUTATION_HPP