Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2025-01-18 09:29:57

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2014 Roshan <thisisroshansmail@gmail.com>
0003 //
0004 // Distributed under the Boost Software License, Version 1.0
0005 // See accompanying file LICENSE_1_0.txt or copy at
0006 // http://www.boost.org/LICENSE_1_0.txt
0007 //
0008 // See http://boostorg.github.com/compute for more information.
0009 //---------------------------------------------------------------------------//
0010 
0011 #ifndef BOOST_COMPUTE_ALGORITHM_NEXT_PERMUTATION_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_NEXT_PERMUTATION_HPP
0013 
0014 #include <iterator>
0015 
0016 #include <boost/static_assert.hpp>
0017 
0018 #include <boost/compute/system.hpp>
0019 #include <boost/compute/command_queue.hpp>
0020 #include <boost/compute/container/detail/scalar.hpp>
0021 #include <boost/compute/algorithm/reverse.hpp>
0022 #include <boost/compute/type_traits/is_device_iterator.hpp>
0023 
0024 namespace boost {
0025 namespace compute {
0026 namespace detail {
0027 
0028 ///
0029 /// \brief Helper function for next_permutation
0030 ///
0031 /// To find rightmost element which is smaller
0032 /// than its next element
0033 ///
0034 template<class InputIterator>
0035 inline InputIterator next_permutation_helper(InputIterator first,
0036                                              InputIterator last,
0037                                              command_queue &queue)
0038 {
0039     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0040 
0041     size_t count = detail::iterator_range_size(first, last);
0042     if(count == 0 || count == 1){
0043         return last;
0044     }
0045     count = count - 1;
0046     const context &context = queue.get_context();
0047 
0048     detail::meta_kernel k("next_permutation");
0049     size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
0050     atomic_max<int_> atomic_max_int;
0051 
0052     k << k.decl<const int_>("i") << " = get_global_id(0);\n"
0053       << k.decl<const value_type>("cur_value") << "="
0054       <<     first[k.var<const int_>("i")] << ";\n"
0055       << k.decl<const value_type>("next_value") << "="
0056       <<     first[k.expr<const int_>("i+1")] << ";\n"
0057       << "if(cur_value < next_value){\n"
0058       << "    " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
0059       << "}\n";
0060 
0061     kernel kernel = k.compile(context);
0062 
0063     scalar<int_> index(context);
0064     kernel.set_arg(index_arg, index.get_buffer());
0065 
0066     index.write(static_cast<int_>(-1), queue);
0067 
0068     queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
0069 
0070     int result = static_cast<int>(index.read(queue));
0071     if(result == -1) return last;
0072     else return first + result;
0073 }
0074 
0075 ///
0076 /// \brief Helper function for next_permutation
0077 ///
0078 /// To find the smallest element to the right of the element found above
0079 /// that is greater than it
0080 ///
0081 template<class InputIterator, class ValueType>
0082 inline InputIterator np_ceiling(InputIterator first,
0083                                 InputIterator last,
0084                                 ValueType value,
0085                                 command_queue &queue)
0086 {
0087     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0088 
0089     size_t count = detail::iterator_range_size(first, last);
0090     if(count == 0){
0091         return last;
0092     }
0093     const context &context = queue.get_context();
0094 
0095     detail::meta_kernel k("np_ceiling");
0096     size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
0097     size_t value_arg = k.add_arg<value_type>(memory_object::private_memory, "value");
0098     atomic_max<int_> atomic_max_int;
0099 
0100     k << k.decl<const int_>("i") << " = get_global_id(0);\n"
0101       << k.decl<const value_type>("cur_value") << "="
0102       <<     first[k.var<const int_>("i")] << ";\n"
0103       << "if(cur_value <= " << first[k.expr<int_>("*index")]
0104       << "      && cur_value > value){\n"
0105       << "    " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
0106       << "}\n";
0107 
0108     kernel kernel = k.compile(context);
0109 
0110     scalar<int_> index(context);
0111     kernel.set_arg(index_arg, index.get_buffer());
0112 
0113     index.write(static_cast<int_>(0), queue);
0114 
0115     kernel.set_arg(value_arg, value);
0116 
0117     queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
0118 
0119     int result = static_cast<int>(index.read(queue));
0120     return first + result;
0121 }
0122 
0123 } // end detail namespace
0124 
0125 ///
0126 /// \brief Permutation generating algorithm
0127 ///
0128 /// Transforms the range [first, last) into the next permutation from the
0129 /// set of all permutations arranged in lexicographic order
0130 /// \return Boolean value signifying if the last permutation was crossed
0131 /// and the range was reset
0132 ///
0133 /// \param first Iterator pointing to start of range
0134 /// \param last Iterator pointing to end of range
0135 /// \param queue Queue on which to execute
0136 ///
0137 /// Space complexity: \Omega(1)
0138 template<class InputIterator>
0139 inline bool next_permutation(InputIterator first,
0140                              InputIterator last,
0141                              command_queue &queue = system::default_queue())
0142 {
0143     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0144     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0145 
0146     if(first == last) return false;
0147 
0148     InputIterator first_element =
0149         detail::next_permutation_helper(first, last, queue);
0150 
0151     if(first_element == last)
0152     {
0153         reverse(first, last, queue);
0154         return false;
0155     }
0156 
0157     value_type first_value = first_element.read(queue);
0158 
0159     InputIterator ceiling_element =
0160         detail::np_ceiling(first_element + 1, last, first_value, queue);
0161 
0162     value_type ceiling_value = ceiling_element.read(queue);
0163 
0164     first_element.write(ceiling_value, queue);
0165     ceiling_element.write(first_value, queue);
0166 
0167     reverse(first_element + 1, last, queue);
0168 
0169     return true;
0170 }
0171 
0172 } // end compute namespace
0173 } // end boost namespace
0174 
0175 #endif // BOOST_COMPUTE_ALGORITHM_NEXT_PERMUTATION_HPP