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_FOR_EACH_HPP
12 #define BOOST_COMPUTE_ALGORITHM_FOR_EACH_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 InputIterator, class Function>
27 struct for_each_kernel : public meta_kernel
28 {
for_each_kernelboost::compute::detail::for_each_kernel29 for_each_kernel(InputIterator first, InputIterator last, Function function)
30 : meta_kernel("for_each")
31 {
32 // store range size
33 m_count = detail::iterator_range_size(first, last);
34
35 // setup kernel source
36 *this << function(first[get_global_id(0)]) << ";\n";
37 }
38
execboost::compute::detail::for_each_kernel39 void exec(command_queue &queue)
40 {
41 exec_1d(queue, 0, m_count);
42 }
43
44 size_t m_count;
45 };
46
47 } // end detail namespace
48
49 /// Calls \p function on each element in the range [\p first, \p last).
50 ///
51 /// Space complexity: \Omega(1)
52 ///
53 /// \see transform()
54 template<class InputIterator, class UnaryFunction>
for_each(InputIterator first,InputIterator last,UnaryFunction function,command_queue & queue=system::default_queue ())55 inline UnaryFunction for_each(InputIterator first,
56 InputIterator last,
57 UnaryFunction function,
58 command_queue &queue = system::default_queue())
59 {
60 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
61
62 detail::for_each_kernel<InputIterator, UnaryFunction> kernel(first, last, function);
63
64 kernel.exec(queue);
65
66 return function;
67 }
68
69 } // end compute namespace
70 } // end boost namespace
71
72 #endif // BOOST_COMPUTE_ALGORITHM_FOR_EACH_HPP
73