Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2025-01-18 09:29:57

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
0003 //
0004 // Distributed under the Boost Software License, Version 1.0
0005 // See accompanying file LICENSE_1_0.txt or copy at
0006 // http://www.boost.org/LICENSE_1_0.txt
0007 //
0008 // See http://boostorg.github.com/compute for more information.
0009 //---------------------------------------------------------------------------//
0010 
0011 #ifndef BOOST_COMPUTE_ALGORITHM_REPLACE_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_REPLACE_HPP
0013 
0014 #include <boost/static_assert.hpp>
0015 
0016 #include <boost/compute/system.hpp>
0017 #include <boost/compute/command_queue.hpp>
0018 #include <boost/compute/detail/meta_kernel.hpp>
0019 #include <boost/compute/detail/iterator_range_size.hpp>
0020 #include <boost/compute/type_traits/is_device_iterator.hpp>
0021 
0022 namespace boost {
0023 namespace compute {
0024 namespace detail {
0025 
0026 template<class Iterator, class T>
0027 class replace_kernel : public meta_kernel
0028 {
0029 public:
0030     replace_kernel()
0031         : meta_kernel("replace")
0032     {
0033         m_count = 0;
0034     }
0035 
0036     void set_range(Iterator first, Iterator last)
0037     {
0038         m_count = detail::iterator_range_size(first, last);
0039 
0040         *this <<
0041             "const uint i = get_global_id(0);\n" <<
0042             "if(" << first[var<cl_uint>("i")] << " == " << var<T>("old_value") << ")\n" <<
0043             "    " << first[var<cl_uint>("i")] << '=' << var<T>("new_value") << ";\n";
0044     }
0045 
0046     void set_old_value(const T &old_value)
0047     {
0048         add_set_arg<T>("old_value", old_value);
0049     }
0050 
0051     void set_new_value(const T &new_value)
0052     {
0053         add_set_arg<T>("new_value", new_value);
0054     }
0055 
0056     void exec(command_queue &queue)
0057     {
0058         if(m_count == 0){
0059             // nothing to do
0060             return;
0061         }
0062 
0063         exec_1d(queue, 0, m_count);
0064     }
0065 
0066 private:
0067     size_t m_count;
0068 };
0069 
0070 } // end detail namespace
0071 
0072 /// Replaces each instance of \p old_value in the range [\p first,
0073 /// \p last) with \p new_value.
0074 ///
0075 /// Space complexity: \Omega(1)
0076 template<class Iterator, class T>
0077 inline void replace(Iterator first,
0078                     Iterator last,
0079                     const T &old_value,
0080                     const T &new_value,
0081                     command_queue &queue = system::default_queue())
0082 {
0083     BOOST_STATIC_ASSERT(is_device_iterator<Iterator>::value);
0084     detail::replace_kernel<Iterator, T> kernel;
0085 
0086     kernel.set_range(first, last);
0087     kernel.set_old_value(old_value);
0088     kernel.set_new_value(new_value);
0089 
0090     kernel.exec(queue);
0091 }
0092 
0093 } // end compute namespace
0094 } // end boost namespace
0095 
0096 #endif // BOOST_COMPUTE_ALGORITHM_REPLACE_HPP