reverse_copy.hpp 2.6 KB

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