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