summaryrefslogtreecommitdiff
path: root/boost/compute/algorithm/scatter.hpp
diff options
context:
space:
mode:
Diffstat (limited to 'boost/compute/algorithm/scatter.hpp')
-rw-r--r--boost/compute/algorithm/scatter.hpp99
1 files changed, 99 insertions, 0 deletions
diff --git a/boost/compute/algorithm/scatter.hpp b/boost/compute/algorithm/scatter.hpp
new file mode 100644
index 0000000000..bea4201628
--- /dev/null
+++ b/boost/compute/algorithm/scatter.hpp
@@ -0,0 +1,99 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@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_HPP
+#define BOOST_COMPUTE_ALGORITHM_SCATTER_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 OutputIterator>
+class scatter_kernel : meta_kernel
+{
+public:
+ scatter_kernel() : meta_kernel("scatter")
+ {}
+
+ void set_range(InputIterator first,
+ InputIterator last,
+ MapIterator map,
+ OutputIterator result)
+ {
+ 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" <<
+ 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.
+///
+/// \see gather()
+template<class InputIterator, class MapIterator, class OutputIterator>
+inline void scatter(InputIterator first,
+ InputIterator last,
+ MapIterator map,
+ OutputIterator result,
+ command_queue &queue = system::default_queue())
+{
+ detail::scatter_kernel<InputIterator, MapIterator, OutputIterator> kernel;
+
+ kernel.set_range(first, last, map, result);
+ kernel.exec(queue);
+}
+
+} // end compute namespace
+} // end boost namespace
+
+#endif // BOOST_COMPUTE_ALGORITHM_SCATTER_HPP