ceph/src/boost/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/static_assert.hpp>
  17
  18 #include <boost/compute/system.hpp>
  19 #include <boost/compute/command_queue.hpp>
  20 #include <boost/compute/algorithm/copy.hpp>
  21 #include <boost/compute/algorithm/reverse.hpp>
  22 #include <boost/compute/type_traits/is_device_iterator.hpp>
  23
  24 namespace boost {
  25 namespace compute {
  26 namespace detail {
  27
  28 template<class Iterator, class OutputIterator>
  29 struct reverse_copy_kernel : public meta_kernel
  30 {
  31     reverse_copy_kernel(Iterator first, Iterator last, OutputIterator result)
  32         : meta_kernel("reverse_copy")
  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             result[var<cl_uint>("j")] << "=" << first[var<cl_uint>("i")] << ";\n";
  42     }
  43
  44     void exec(command_queue &queue)
  45     {
  46         exec_1d(queue, 0, m_size);
  47     }
  48
  49     size_t m_size;
  50 };
  51
  52 } // end detail namespace
  53
  54 /// Copies the elements in the range [\p first, \p last) in reversed
  55 /// order to the range beginning at \p result.
  56 ///
  57 /// Space complexity: \Omega(1)
  58 ///
  59 /// \see reverse()
  60 template<class InputIterator, class OutputIterator>
  61 inline OutputIterator
  62 reverse_copy(InputIterator first,
  63              InputIterator last,
  64              OutputIterator result,
  65              command_queue &queue = system::default_queue())
  66 {
  67     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
  68     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
  69
  70     typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
  71
  72     difference_type count = std::distance(first, last);
  73
  74     detail::reverse_copy_kernel<InputIterator, OutputIterator>
  75         kernel(first, last, result);
  76
  77     // run kernel
  78     kernel.exec(queue);
  79
  80     // return iterator to the end of result
  81     return result + count;
  82 }
  83
  84 } // end compute namespace
  85 } // end boost namespace
  86
  87 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_COPY_HPP