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