compute/algorithm/reverse_copy.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_COPY_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_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/algorithm/copy.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 template<class Iterator, class OutputIterator>
0029 struct reverse_copy_kernel : public meta_kernel
0030 {
0031     reverse_copy_kernel(Iterator first, Iterator last, OutputIterator result)
0032         : meta_kernel("reverse_copy")
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             result[var<cl_uint>("j")] << "=" << first[var<cl_uint>("i")] << ";\n";
0042     }
0043
0044     void exec(command_queue &queue)
0045     {
0046         exec_1d(queue, 0, m_size);
0047     }
0048
0049     size_t m_size;
0050 };
0051
0052 } // end detail namespace
0053
0054 /// Copies the elements in the range [\p first, \p last) in reversed
0055 /// order to the range beginning at \p result.
0056 ///
0057 /// Space complexity: \Omega(1)
0058 ///
0059 /// \see reverse()
0060 template<class InputIterator, class OutputIterator>
0061 inline OutputIterator
0062 reverse_copy(InputIterator first,
0063              InputIterator last,
0064              OutputIterator result,
0065              command_queue &queue = system::default_queue())
0066 {
0067     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0068     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
0069
0070     typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
0071
0072     difference_type count = std::distance(first, last);
0073
0074     detail::reverse_copy_kernel<InputIterator, OutputIterator>
0075         kernel(first, last, result);
0076
0077     // run kernel
0078     kernel.exec(queue);
0079
0080     // return iterator to the end of result
0081     return result + count;
0082 }
0083
0084 } // end compute namespace
0085 } // end boost namespace
0086
0087 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP