• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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:
replace_kernel()30     replace_kernel()
31         : meta_kernel("replace")
32     {
33         m_count = 0;
34     }
35 
set_range(Iterator first,Iterator last)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 
set_old_value(const T & old_value)46     void set_old_value(const T &old_value)
47     {
48         add_set_arg<T>("old_value", old_value);
49     }
50 
set_new_value(const T & new_value)51     void set_new_value(const T &new_value)
52     {
53         add_set_arg<T>("new_value", new_value);
54     }
55 
exec(command_queue & queue)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>
replace(Iterator first,Iterator last,const T & old_value,const T & new_value,command_queue & queue=system::default_queue ())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
97