diff options
Diffstat (limited to 'boost/compute/algorithm/scatter_if.hpp')
-rw-r--r-- | boost/compute/algorithm/scatter_if.hpp | 119 |
1 files changed, 119 insertions, 0 deletions
diff --git a/boost/compute/algorithm/scatter_if.hpp b/boost/compute/algorithm/scatter_if.hpp new file mode 100644 index 0000000000..159edd8c86 --- /dev/null +++ b/boost/compute/algorithm/scatter_if.hpp @@ -0,0 +1,119 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2015 Jakub Pola <jakub.pola@gmail.com> +// +// Distributed under the Boost Software License, Version 1.0 +// See accompanying file LICENSE_1_0.txt or copy at +// http://www.boost.org/LICENSE_1_0.txt +// +// See http://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP +#define BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP + +#include <boost/algorithm/string/replace.hpp> + +#include <boost/compute/system.hpp> +#include <boost/compute/exception.hpp> +#include <boost/compute/command_queue.hpp> +#include <boost/compute/iterator/buffer_iterator.hpp> +#include <boost/compute/type_traits/type_name.hpp> +#include <boost/compute/detail/iterator_range_size.hpp> +#include <boost/compute/detail/meta_kernel.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator, class Predicate> +class scatter_if_kernel : meta_kernel +{ +public: + scatter_if_kernel() : meta_kernel("scatter_if") + {} + + void set_range(InputIterator first, + InputIterator last, + MapIterator map, + StencilIterator stencil, + OutputIterator result, + Predicate predicate) + { + m_count = iterator_range_size(first, last); + m_input_offset = first.get_index(); + m_output_offset = result.get_index(); + + m_input_offset_arg = add_arg<uint_>("input_offset"); + m_output_offset_arg = add_arg<uint_>("output_offset"); + + *this << + "const uint i = get_global_id(0);\n" << + "uint i1 = " << map[expr<uint_>("i")] << + " + output_offset;\n" << + "uint i2 = i + input_offset;\n" << + if_(predicate(stencil[expr<uint_>("i")])) << "\n" << + result[expr<uint_>("i1")] << "=" << + first[expr<uint_>("i2")] << ";\n"; + } + + event exec(command_queue &queue) + { + if(m_count == 0) { + return event(); + } + + set_arg(m_input_offset_arg, uint_(m_input_offset)); + set_arg(m_output_offset_arg, uint_(m_output_offset)); + + return exec_1d(queue, 0, m_count); + } + +private: + size_t m_count; + size_t m_input_offset; + size_t m_input_offset_arg; + size_t m_output_offset; + size_t m_output_offset_arg; +}; + +} // end detail namespace + +/// Copies the elements from the range [\p first, \p last) to the range +/// beginning at \p result using the output indices from the range beginning +/// at \p map if stencil is resolved to true. By default the predicate is +/// an identity +/// +/// +template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator, + class Predicate> +inline void scatter_if(InputIterator first, + InputIterator last, + MapIterator map, + StencilIterator stencil, + OutputIterator result, + Predicate predicate, + command_queue &queue = system::default_queue()) +{ + detail::scatter_if_kernel<InputIterator, MapIterator, StencilIterator, OutputIterator, Predicate> kernel; + + kernel.set_range(first, last, map, stencil, result, predicate); + kernel.exec(queue); +} + +template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator> +inline void scatter_if(InputIterator first, + InputIterator last, + MapIterator map, + StencilIterator stencil, + OutputIterator result, + command_queue &queue = system::default_queue()) +{ + typedef typename std::iterator_traits<StencilIterator>::value_type T; + + scatter_if(first, last, map, stencil, result, identity<T>(), queue); +} + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP |