ceph/src/boost/boost/compute/algorithm/replace.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_REPLACE_HPP
  12 #define BOOST_COMPUTE_ALGORITHM_REPLACE_HPP
  13
  14 #include <boost/static_assert.hpp>
  15
  16 #include <boost/compute/system.hpp>
  17 #include <boost/compute/command_queue.hpp>
  18 #include <boost/compute/detail/meta_kernel.hpp>
  19 #include <boost/compute/detail/iterator_range_size.hpp>
  20 #include <boost/compute/type_traits/is_device_iterator.hpp>
  21
  22 namespace boost {
  23 namespace compute {
  24 namespace detail {
  25
  26 template<class Iterator, class T>
  27 class replace_kernel : public meta_kernel
  28 {
  29 public:
  30     replace_kernel()
  31         : meta_kernel("replace")
  32     {
  33         m_count = 0;
  34     }
  35
  36     void set_range(Iterator first, Iterator last)
  37     {
  38         m_count = detail::iterator_range_size(first, last);
  39
  40         *this <<
  41             "const uint i = get_global_id(0);\n" <<
  42             "if(" << first[var<cl_uint>("i")] << " == " << var<T>("old_value") << ")\n" <<
  43             "    " << first[var<cl_uint>("i")] << '=' << var<T>("new_value") << ";\n";
  44     }
  45
  46     void set_old_value(const T &old_value)
  47     {
  48         add_set_arg<T>("old_value", old_value);
  49     }
  50
  51     void set_new_value(const T &new_value)
  52     {
  53         add_set_arg<T>("new_value", new_value);
  54     }
  55
  56     void exec(command_queue &queue)
  57     {
  58         if(m_count == 0){
  59             // nothing to do
  60             return;
  61         }
  62
  63         exec_1d(queue, 0, m_count);
  64     }
  65
  66 private:
  67     size_t m_count;
  68 };
  69
  70 } // end detail namespace
  71
  72 /// Replaces each instance of \p old_value in the range [\p first,
  73 /// \p last) with \p new_value.
  74 ///
  75 /// Space complexity: \Omega(1)
  76 template<class Iterator, class T>
  77 inline void replace(Iterator first,
  78                     Iterator last,
  79                     const T &old_value,
  80                     const T &new_value,
  81                     command_queue &queue = system::default_queue())
  82 {
  83     BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
  84     detail::replace_kernel<Iterator, T> kernel;
  85
  86     kernel.set_range(first, last);
  87     kernel.set_old_value(old_value);
  88     kernel.set_new_value(new_value);
  89
  90     kernel.exec(queue);
  91 }
  92
  93 } // end compute namespace
  94 } // end boost namespace
  95
  96 #endif // BOOST_COMPUTE_ALGORITHM_REPLACE_HPP