compute/algorithm/reverse.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@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_REVERSE_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
0013
0014 #include <boost/static_assert.hpp>
0015
0016 #include <boost/compute/system.hpp>
0017 #include <boost/compute/command_queue.hpp>
0018 #include <boost/compute/detail/meta_kernel.hpp>
0019 #include <boost/compute/detail/iterator_range_size.hpp>
0020 #include <boost/compute/type_traits/is_device_iterator.hpp>
0021
0022 namespace boost {
0023 namespace compute {
0024 namespace detail {
0025
0026 template<class Iterator>
0027 struct reverse_kernel : public meta_kernel
0028 {
0029     reverse_kernel(Iterator first, Iterator last)
0030         : meta_kernel("reverse")
0031     {
0032         typedef typename std::iterator_traits<Iterator>::value_type value_type;
0033
0034         // store size of the range
0035         m_size = detail::iterator_range_size(first, last);
0036         add_set_arg<const cl_uint>("size", static_cast<const cl_uint>(m_size));
0037
0038         *this <<
0039             decl<cl_uint>("i") << " = get_global_id(0);\n" <<
0040             decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" <<
0041             decl<value_type>("tmp") << "=" << first[var<cl_uint>("i")] << ";\n" <<
0042             first[var<cl_uint>("i")] << "=" << first[var<cl_uint>("j")] << ";\n" <<
0043             first[var<cl_uint>("j")] << "= tmp;\n";
0044     }
0045
0046     void exec(command_queue &queue)
0047     {
0048         exec_1d(queue, 0, m_size / 2);
0049     }
0050
0051     size_t m_size;
0052 };
0053
0054 } // end detail namespace
0055
0056 /// Reverses the elements in the range [\p first, \p last).
0057 ///
0058 /// Space complexity: \Omega(1)
0059 ///
0060 /// \see reverse_copy()
0061 template<class Iterator>
0062 inline void reverse(Iterator first,
0063                     Iterator last,
0064                     command_queue &queue = system::default_queue())
0065 {
0066     BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
0067     size_t count = detail::iterator_range_size(first, last);
0068     if(count < 2){
0069         return;
0070     }
0071
0072     detail::reverse_kernel<Iterator> kernel(first, last);
0073
0074     kernel.exec(queue);
0075 }
0076
0077 } // end compute namespace
0078 } // end boost namespace
0079
0080 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_HPP