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