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