ceph/src/boost/boost/compute/algorithm/gather.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_GATHER_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_GATHER_HPP
  13
  14 #include <boost/static_assert.hpp>
  15
  16 #include <boost/compute/command_queue.hpp>
  17 #include <boost/compute/detail/iterator_range_size.hpp>
  18 #include <boost/compute/detail/meta_kernel.hpp>
  19 #include <boost/compute/exception.hpp>
  20 #include <boost/compute/iterator/buffer_iterator.hpp>
  21 #include <boost/compute/system.hpp>
  22 #include <boost/compute/type_traits/type_name.hpp>
  23 #include <boost/compute/type_traits/is_device_iterator.hpp>
  24
  25 namespace boost {
  26 namespace compute {
  27 namespace detail {
  28
  29 template<class InputIterator, class MapIterator, class OutputIterator>
  30 class gather_kernel : public meta_kernel
  31 {
  32 public:
  33     gather_kernel() : meta_kernel("gather")
  34     {}
  35
  36     void set_range(MapIterator first,
  37                    MapIterator last,
  38                    InputIterator input,
  39                    OutputIterator result)
  40     {
  41         m_count = iterator_range_size(first, last);
  42
  43         *this <<
  44             "const uint i = get_global_id(0);\n" <<
  45             result[expr<uint_>("i")] << "=" <<
  46                 input[first[expr<uint_>("i")]] << ";\n";
  47     }
  48
  49     event exec(command_queue &queue)
  50     {
  51         if(m_count == 0) {
  52             return event();
  53         }
  54
  55         return exec_1d(queue, 0, m_count);
  56     }
  57
  58 private:
  59     size_t m_count;
  60 };
  61
  62 } // end detail namespace
  63
  64 /// Copies the elements using the indices from the range [\p first, \p last)
  65 /// to the range beginning at \p result using the input values from the range
  66 /// beginning at \p input.
  67 ///
  68 /// Space complexity: \Omega(1)
  69 ///
  70 /// \see scatter()
  71 template<class InputIterator, class MapIterator, class OutputIterator>
  72 inline void gather(MapIterator first,
  73                    MapIterator last,
  74                    InputIterator input,
  75                    OutputIterator result,
  76                    command_queue &queue = system::default_queue())
  77 {
  78     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
  79     BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
  80     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
  81
  82     detail::gather_kernel<InputIterator, MapIterator, OutputIterator> kernel;
  83
  84     kernel.set_range(first, last, input, result);
  85     kernel.exec(queue);
  86 }
  87
  88 } // end compute namespace
  89 } // end boost namespace
  90
  91 #endif // BOOST_COMPUTE_ALGORITHM_GATHER_HPP