#ifndef BOOST_COMPUTE_ALGORITHM_FOR_EACH_HPP
#define BOOST_COMPUTE_ALGORITHM_FOR_EACH_HPP
#include <boost/static_assert.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/type_traits/is_device_iterator.hpp>
namespace boost {
namespace compute {
namespace detail {
template<class InputIterator, class Function>
struct for_each_kernel : public meta_kernel
{
for_each_kernel(InputIterator first, InputIterator last, Function function)
: meta_kernel("for_each")
{
m_count = detail::iterator_range_size(first, last);
*this << function(first[get_global_id(0)]) << ";\n";
}
void exec(command_queue &queue)
{
exec_1d(queue, 0, m_count);
}
size_t m_count;
};
}
template<class InputIterator, class UnaryFunction>
inline UnaryFunction for_each(InputIterator first,
InputIterator last,
UnaryFunction function,
command_queue &queue = system::default_queue())
{
BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
detail::for_each_kernel<InputIterator, UnaryFunction> kernel(first, last, function);
kernel.exec(queue);
return function;
}
} }
#endif