projects / ceph.git / blame
| Commit | Line | Data |
|---|---|---|
| 7c673cae FG |
1 | //---------------------------------------------------------------------------// |
| 2 | // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@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_REVERSE_HPP | |
| 12 | #define BOOST_COMPUTE_ALGORITHM_REVERSE_HPP | |
| 13 | ||
| 14 | #include <boost/compute/system.hpp> | |
| 15 | #include <boost/compute/command_queue.hpp> | |
| 16 | #include <boost/compute/detail/meta_kernel.hpp> | |
| 17 | #include <boost/compute/detail/iterator_range_size.hpp> | |
| 18 | ||
| 19 | namespace boost { | |
| 20 | namespace compute { | |
| 21 | namespace detail { | |
| 22 | ||
| 23 | template<class Iterator> | |
| 24 | struct reverse_kernel : public meta_kernel | |
| 25 | { | |
| 26 | reverse_kernel(Iterator first, Iterator last) | |
| 27 | : meta_kernel("reverse") | |
| 28 | { | |
| 29 | typedef typename std::iterator_traits<Iterator>::value_type value_type; | |
| 30 | ||
| 31 | // store size of the range | |
| 32 | m_size = detail::iterator_range_size(first, last); | |
| 33 | add_set_arg<const cl_uint>("size", static_cast<const cl_uint>(m_size)); | |
| 34 | ||
| 35 | *this << | |
| 36 | decl<cl_uint>("i") << " = get_global_id(0);\n" << | |
| 37 | decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" << | |
| 38 | decl<value_type>("tmp") << "=" << first[var<cl_uint>("i")] << ";\n" << | |
| 39 | first[var<cl_uint>("i")] << "=" << first[var<cl_uint>("j")] << ";\n" << | |
| 40 | first[var<cl_uint>("j")] << "= tmp;\n"; | |
| 41 | } | |
| 42 | ||
| 43 | void exec(command_queue &queue) | |
| 44 | { | |
| 45 | exec_1d(queue, 0, m_size / 2); | |
| 46 | } | |
| 47 | ||
| 48 | size_t m_size; | |
| 49 | }; | |
| 50 | ||
| 51 | } // end detail namespace | |
| 52 | ||
| 53 | /// Reverses the elements in the range [\p first, \p last). | |
| 54 | /// | |
| 55 | /// \see reverse_copy() | |
| 56 | template<class Iterator> | |
| 57 | inline void reverse(Iterator first, | |
| 58 | Iterator last, | |
| 59 | command_queue &queue = system::default_queue()) | |
| 60 | { | |
| 61 | size_t count = detail::iterator_range_size(first, last); | |
| 62 | if(count < 2){ | |
| 63 | return; | |
| 64 | } | |
| 65 | ||
| 66 | detail::reverse_kernel<Iterator> kernel(first, last); | |
| 67 | ||
| 68 | kernel.exec(queue); | |
| 69 | } | |
| 70 | ||
| 71 | } // end compute namespace | |
| 72 | } // end boost namespace | |
| 73 | ||
| 74 | #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_HPP |