ceph/src/boost/libs/compute/include/boost/compute/algorithm/reverse_copy.hpp

   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_COPY_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP
  13
  14 #include <iterator>
  15
  16 #include <boost/compute/system.hpp>
  17 #include <boost/compute/command_queue.hpp>
  18 #include <boost/compute/algorithm/copy.hpp>
  19 #include <boost/compute/algorithm/reverse.hpp>
  20
  21 namespace boost {
  22 namespace compute {
  23 namespace detail {
  24
  25 template<class Iterator, class OutputIterator>
  26 struct reverse_copy_kernel : public meta_kernel
  27 {
  28     reverse_copy_kernel(Iterator first, Iterator last, OutputIterator result)
  29         : meta_kernel("reverse_copy")
  30     {
  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));
  34
  35         *this <<
  36             decl<cl_uint>("i") << " = get_global_id(0);\n" <<
  37             decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" <<
  38             result[var<cl_uint>("j")] << "=" << first[var<cl_uint>("i")] << ";\n";
  39     }
  40
  41     void exec(command_queue &queue)
  42     {
  43         exec_1d(queue, 0, m_size);
  44     }
  45
  46     size_t m_size;
  47 };
  48
  49 } // end detail namespace
  50
  51 /// Copies the elements in the range [\p first, \p last) in reversed
  52 /// order to the range beginning at \p result.
  53 ///
  54 /// \see reverse()
  55 template<class InputIterator, class OutputIterator>
  56 inline OutputIterator
  57 reverse_copy(InputIterator first,
  58              InputIterator last,
  59              OutputIterator result,
  60              command_queue &queue = system::default_queue())
  61 {
  62     typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
  63
  64     difference_type count = std::distance(first, last);
  65
  66     detail::reverse_copy_kernel<InputIterator, OutputIterator>
  67         kernel(first, last, result);
  68
  69     // run kernel
  70     kernel.exec(queue);
  71
  72     // return iterator to the end of result
  73     return result + count;
  74 }
  75
  76 } // end compute namespace
  77 } // end boost namespace
  78
  79 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP