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/static_assert.hpp>
15
16 #include <boost/compute/system.hpp>
17 #include <boost/compute/command_queue.hpp>
18 #include <boost/compute/detail/meta_kernel.hpp>
19 #include <boost/compute/detail/iterator_range_size.hpp>
20 #include <boost/compute/type_traits/is_device_iterator.hpp>
21
22 namespace boost {
23 namespace compute {
24 namespace detail {
25
26 template<class Iterator>
27 struct reverse_kernel : public meta_kernel
28 {
reverse_kernelboost::compute::detail::reverse_kernel29 reverse_kernel(Iterator first, Iterator last)
30 : meta_kernel("reverse")
31 {
32 typedef typename std::iterator_traits<Iterator>::value_type value_type;
33
34 // store size of the range
35 m_size = detail::iterator_range_size(first, last);
36 add_set_arg<const cl_uint>("size", static_cast<const cl_uint>(m_size));
37
38 *this <<
39 decl<cl_uint>("i") << " = get_global_id(0);\n" <<
40 decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" <<
41 decl<value_type>("tmp") << "=" << first[var<cl_uint>("i")] << ";\n" <<
42 first[var<cl_uint>("i")] << "=" << first[var<cl_uint>("j")] << ";\n" <<
43 first[var<cl_uint>("j")] << "= tmp;\n";
44 }
45
execboost::compute::detail::reverse_kernel46 void exec(command_queue &queue)
47 {
48 exec_1d(queue, 0, m_size / 2);
49 }
50
51 size_t m_size;
52 };
53
54 } // end detail namespace
55
56 /// Reverses the elements in the range [\p first, \p last).
57 ///
58 /// Space complexity: \Omega(1)
59 ///
60 /// \see reverse_copy()
61 template<class Iterator>
reverse(Iterator first,Iterator last,command_queue & queue=system::default_queue ())62 inline void reverse(Iterator first,
63 Iterator last,
64 command_queue &queue = system::default_queue())
65 {
66 BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
67 size_t count = detail::iterator_range_size(first, last);
68 if(count < 2){
69 return;
70 }
71
72 detail::reverse_kernel<Iterator> kernel(first, last);
73
74 kernel.exec(queue);
75 }
76
77 } // end compute namespace
78 } // end boost namespace
79
80 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
81