[compute] 08/46: scatter if operation
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Mon Dec 21 18:28:37 UTC 2015
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch master
in repository compute.
commit f25b47d6f304fd9a90f6ee30db5250da912d126f
Author: Jakub Pola <jakub.pola at gmail.com>
Date: Mon Sep 28 17:17:10 2015 +0200
scatter if operation
---
include/boost/compute/algorithm/scatter_if.hpp | 119 +++++++++++++++++++++
test/CMakeLists.txt | 1 +
test/test_scatter_if.cpp | 141 +++++++++++++++++++++++++
3 files changed, 261 insertions(+)
diff --git a/include/boost/compute/algorithm/scatter_if.hpp b/include/boost/compute/algorithm/scatter_if.hpp
new file mode 100644
index 0000000..766cec4
--- /dev/null
+++ b/include/boost/compute/algorithm/scatter_if.hpp
@@ -0,0 +1,119 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2013 Kyle Lutz <kyle.r.lutz at 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
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index be12385..d391682 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -117,6 +117,7 @@ add_compute_test("algorithm.rotate" test_rotate.cpp)
add_compute_test("algorithm.rotate_copy" test_rotate_copy.cpp)
add_compute_test("algorithm.scan" test_scan.cpp)
add_compute_test("algorithm.scatter" test_scatter.cpp)
+add_compute_test("algorithm.scatter_if" test_scatter_if.cpp)
add_compute_test("algorithm.search" test_search.cpp)
add_compute_test("algorithm.search_n" test_search_n.cpp)
add_compute_test("algorithm.set_difference" test_set_difference.cpp)
diff --git a/test/test_scatter_if.cpp b/test/test_scatter_if.cpp
new file mode 100644
index 0000000..d564449
--- /dev/null
+++ b/test/test_scatter_if.cpp
@@ -0,0 +1,141 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2013 Kyle Lutz <kyle.r.lutz at 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.
+//---------------------------------------------------------------------------//
+
+#define BOOST_TEST_MODULE TestScatter
+#include <boost/test/unit_test.hpp>
+
+#include <boost/compute/system.hpp>
+#include <boost/compute/algorithm/scatter.hpp>
+#include <boost/compute/algorithm/scatter_if.hpp>
+#include <boost/compute/container/vector.hpp>
+#include <boost/compute/iterator/constant_buffer_iterator.hpp>
+#include <boost/compute/iterator/counting_iterator.hpp>
+#include <boost/compute/functional.hpp>
+#include "check_macros.hpp"
+#include "context_setup.hpp"
+
+namespace bc = boost::compute;
+
+BOOST_AUTO_TEST_CASE(scatter_if_int)
+{
+ int input_data[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
+ bc::vector<int> input(input_data, input_data + 10, queue);
+
+ int map_data[] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
+
+ bc::vector<int> map(map_data, map_data + 10, queue);
+
+ int stencil_data[] = {0, 1, 0, 1, 0, 1, 0, 1, 0, 1};
+
+ bc::vector<uint> stencil(stencil_data, stencil_data + 10, queue);
+
+ bc::vector<int> output = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
+
+ bc::scatter_if(input.begin(), input.end(),
+ map.begin(), stencil.begin(),
+ output.begin());
+
+ CHECK_RANGE_EQUAL(int, 10, output, (9, -1, 7, -1, 5, -1, 3, -1, 1, -1) );
+}
+
+BOOST_AUTO_TEST_CASE(scatter_if_constant_indices)
+{
+ int input_data[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
+ bc::vector<int> input(input_data, input_data + 10, queue);
+
+ int map_data[] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
+ bc::buffer map_buffer(context,
+ 10 * sizeof(int),
+ bc::buffer::read_only | bc::buffer::use_host_ptr,
+ map_data);
+
+ int stencil_data[] = {0, 1, 0, 1, 0, 1, 0, 1, 0, 1};
+ bc::buffer stencil_buffer(context,
+ 10 * sizeof(int),
+ bc::buffer::read_only | bc::buffer::use_host_ptr,
+ stencil_data);
+
+ bc::vector<int> output = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
+
+ bc::scatter_if(input.begin(),
+ input.end(),
+ bc::make_constant_buffer_iterator<int>(map_buffer, 0),
+ bc::make_constant_buffer_iterator<int>(stencil_buffer, 0),
+ output.begin(),
+ queue);
+
+ CHECK_RANGE_EQUAL(int, 10, output, (9, -1, 7, -1, 5, -1, 3, -1, 1, -1) );
+}
+
+
+BOOST_AUTO_TEST_CASE(scatter_if_function)
+{
+ int input_data[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
+ bc::vector<int> input(input_data, input_data + 10, queue);
+
+ int map_data[] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
+ bc::vector<int> map(map_data, map_data + 10, queue);
+
+ int stencil_data[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
+ bc::vector<uint> stencil(stencil_data, stencil_data + 10, queue);
+
+ bc::vector<int> output = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
+
+ BOOST_COMPUTE_FUNCTION(int, gt_than_5, (int x),
+ {
+ if (x > 5)
+ return true;
+ else
+ return false;
+ });
+
+ bc::scatter_if(input.begin(),
+ input.end(),
+ map.begin(),
+ stencil.begin(),
+ output.begin(),
+ gt_than_5,
+ queue);
+
+ CHECK_RANGE_EQUAL(int, 10, output, (9, 8, 7, 6, -1, -1, -1, -1, -1, -1) );
+}
+
+
+BOOST_AUTO_TEST_CASE(scatter_if_counting_iterator)
+{
+ int input_data[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
+ bc::vector<int> input(input_data, input_data + 10, queue);
+
+ int map_data[] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
+ bc::vector<int> map(map_data, map_data + 10, queue);
+
+ bc::vector<int> output = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
+
+ BOOST_COMPUTE_FUNCTION(int, gt_than_5, (int x),
+ {
+ if (x > 5)
+ return true;
+ else
+ return false;
+ });
+
+ bc::scatter_if(input.begin(),
+ input.end(),
+ map.begin(),
+ bc::make_counting_iterator<int>(0),
+ output.begin(),
+ gt_than_5,
+ queue);
+
+ CHECK_RANGE_EQUAL(int, 10, output, (9, 8, 7, 6, -1, -1, -1, -1, -1, -1) );
+
+}
+
+BOOST_AUTO_TEST_SUITE_END()
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/compute.git
More information about the debian-science-commits
mailing list