1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
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
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
11 #ifndef BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
12 #define BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
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>
23 template<class Iterator>
24 struct reverse_kernel : public meta_kernel
26 reverse_kernel(Iterator first, Iterator last)
27 : meta_kernel("reverse")
29 typedef typename std::iterator_traits<Iterator>::value_type value_type;
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));
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";
43 void exec(command_queue &queue)
45 exec_1d(queue, 0, m_size / 2);
51 } // end detail namespace
53 /// Reverses the elements in the range [\p first, \p last).
55 /// \see reverse_copy()
56 template<class Iterator>
57 inline void reverse(Iterator first,
59 command_queue &queue = system::default_queue())
61 size_t count = detail::iterator_range_size(first, last);
66 detail::reverse_kernel<Iterator> kernel(first, last);
71 } // end compute namespace
72 } // end boost namespace
74 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_HPP