123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687 |
- //---------------------------------------------------------------------------//
- // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
- //
- // Distributed under the Boost Software License, Version 1.0
- // See accompanying file LICENSE_1_0.txt or copy at
- // http://www.boost.org/LICENSE_1_0.txt
- //
- // See http://boostorg.github.com/compute for more information.
- //---------------------------------------------------------------------------//
- #ifndef BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP
- #define BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP
- #include <iterator>
- #include <boost/static_assert.hpp>
- #include <boost/compute/system.hpp>
- #include <boost/compute/command_queue.hpp>
- #include <boost/compute/algorithm/copy.hpp>
- #include <boost/compute/algorithm/reverse.hpp>
- #include <boost/compute/type_traits/is_device_iterator.hpp>
- namespace boost {
- namespace compute {
- namespace detail {
- template<class Iterator, class OutputIterator>
- struct reverse_copy_kernel : public meta_kernel
- {
- reverse_copy_kernel(Iterator first, Iterator last, OutputIterator result)
- : meta_kernel("reverse_copy")
- {
- // store size of the range
- m_size = detail::iterator_range_size(first, last);
- add_set_arg<const cl_uint>("size", static_cast<const cl_uint>(m_size));
- *this <<
- decl<cl_uint>("i") << " = get_global_id(0);\n" <<
- decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" <<
- result[var<cl_uint>("j")] << "=" << first[var<cl_uint>("i")] << ";\n";
- }
- void exec(command_queue &queue)
- {
- exec_1d(queue, 0, m_size);
- }
- size_t m_size;
- };
- } // end detail namespace
- /// Copies the elements in the range [\p first, \p last) in reversed
- /// order to the range beginning at \p result.
- ///
- /// Space complexity: \Omega(1)
- ///
- /// \see reverse()
- template<class InputIterator, class OutputIterator>
- inline OutputIterator
- reverse_copy(InputIterator first,
- InputIterator last,
- OutputIterator result,
- command_queue &queue = system::default_queue())
- {
- BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
- BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
- typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
- difference_type count = std::distance(first, last);
- detail::reverse_copy_kernel<InputIterator, OutputIterator>
- kernel(first, last, result);
- // run kernel
- kernel.exec(queue);
- // return iterator to the end of result
- return result + count;
- }
- } // end compute namespace
- } // end boost namespace
- #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP
|