ceph/src/boost/boost/compute/algorithm/reverse.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_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_REVERSE_HPP
  13
  14 #include <boost/compute/system.hpp>
  15 #include <boost/compute/command_queue.hpp>
  16 #include <boost/compute/detail/meta_kernel.hpp>
  17 #include <boost/compute/detail/iterator_range_size.hpp>
  18
  19 namespace boost {
  20 namespace compute {
  21 namespace detail {
  22
  23 template<class Iterator>
  24 struct reverse_kernel : public meta_kernel
  25 {
  26     reverse_kernel(Iterator first, Iterator last)
  27         : meta_kernel("reverse")
  28     {
  29         typedef typename std::iterator_traits<Iterator>::value_type value_type;
  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             decl<value_type>("tmp") << "=" << first[var<cl_uint>("i")] << ";\n" <<
  39             first[var<cl_uint>("i")] << "=" << first[var<cl_uint>("j")] << ";\n" <<
  40             first[var<cl_uint>("j")] << "= tmp;\n";
  41     }
  42
  43     void exec(command_queue &queue)
  44     {
  45         exec_1d(queue, 0, m_size / 2);
  46     }
  47
  48     size_t m_size;
  49 };
  50
  51 } // end detail namespace
  52
  53 /// Reverses the elements in the range [\p first, \p last).
  54 ///
  55 /// Space complexity: \Omega(1)
  56 ///
  57 /// \see reverse_copy()
  58 template<class Iterator>
  59 inline void reverse(Iterator first,
  60                     Iterator last,
  61                     command_queue &queue = system::default_queue())
  62 {
  63     size_t count = detail::iterator_range_size(first, last);
  64     if(count < 2){
  65         return;
  66     }
  67
  68     detail::reverse_kernel<Iterator> kernel(first, last);
  69
  70     kernel.exec(queue);
  71 }
  72
  73 } // end compute namespace
  74 } // end boost namespace
  75
  76 #endif // BOOST_COMPUTE_ALGORITHM_REVERSE_HPP