]> git.proxmox.com Git - ceph.git/blame - ceph/src/boost/libs/compute/include/boost/compute/algorithm/next_permutation.hpp
bump version to 12.2.2-pve1
[ceph.git] / ceph / src / boost / libs / compute / include / boost / compute / algorithm / next_permutation.hpp
CommitLineData
7c673cae
FG
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
21namespace boost {
22namespace compute {
23namespace 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///
31template<class InputIterator>
32inline 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///
78template<class InputIterator, class ValueType>
79inline 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///
134template<class InputIterator>
135inline 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