/usr/include/boost/compute/algorithm
Edit: /usr/include/boost/compute/algorithm/prev_permutation.hpp (5695B)
//---------------------------------------------------------------------------//
// Copyright (c) 2014 Roshan
//
// 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_PREV_PERMUTATION_HPP
#define BOOST_COMPUTE_ALGORITHM_PREV_PERMUTATION_HPP
#include
#include
#include
#include
#include
#include
#include
namespace boost {
namespace compute {
namespace detail {
///
/// \brief Helper function for prev_permutation
///
/// To find rightmost element which is greater
/// than its next element
///
template
inline InputIterator prev_permutation_helper(InputIterator first,
InputIterator last,
command_queue &queue)
{
typedef typename std::iterator_traits::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("prev_permutation");
size_t index_arg = k.add_arg(memory_object::global_memory, "index");
atomic_max atomic_max_int;
k << k.decl("i") << " = get_global_id(0);\n"
<< k.decl("cur_value") << "="
<< first[k.var("i")] << ";\n"
<< k.decl("next_value") << "="
<< first[k.expr("i+1")] << ";\n"
<< "if(cur_value > next_value){\n"
<< " " << atomic_max_int(k.var("index"), k.var("i")) << ";\n"
<< "}\n";
kernel kernel = k.compile(context);
scalar index(context);
kernel.set_arg(index_arg, index.get_buffer());
index.write(static_cast(-1), queue);
queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
int result = static_cast(index.read(queue));
if(result == -1) return last;
else return first + result;
}
///
/// \brief Helper function for prev_permutation
///
/// To find the largest element to the right of the element found above
/// that is smaller than it
///
template
inline InputIterator pp_floor(InputIterator first,
InputIterator last,
ValueType value,
command_queue &queue)
{
typedef typename std::iterator_traits::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("pp_floor");
size_t index_arg = k.add_arg(memory_object::global_memory, "index");
size_t value_arg = k.add_arg(memory_object::private_memory, "value");
atomic_max atomic_max_int;
k << k.decl("i") << " = get_global_id(0);\n"
<< k.decl("cur_value") << "="
<< first[k.var("i")] << ";\n"
<< "if(cur_value >= " << first[k.expr("*index")]
<< " && cur_value < value){\n"
<< " " << atomic_max_int(k.var("index"), k.var("i")) << ";\n"
<< "}\n";
kernel kernel = k.compile(context);
scalar index(context);
kernel.set_arg(index_arg, index.get_buffer());
index.write(static_cast(0), queue);
kernel.set_arg(value_arg, value);
queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
int result = static_cast(index.read(queue));
return first + result;
}
} // end detail namespace
///
/// \brief Permutation generating algorithm
///
/// Transforms the range [first, last) into the previous permutation from
/// the set of all permutations arranged in lexicographic order
/// \return Boolean value signifying if the first 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
inline bool prev_permutation(InputIterator first,
InputIterator last,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator::value);
typedef typename std::iterator_traits::value_type value_type;
if(first == last) return false;
InputIterator first_element =
detail::prev_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::pp_floor(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_PREV_PERMUTATION_HPP