1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2015 Jakub Pola <jakub.pola@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_SCATTER_IF_HPP
12 #define BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP
13
14 #include <boost/static_assert.hpp>
15 #include <boost/algorithm/string/replace.hpp>
16
17 #include <boost/compute/system.hpp>
18 #include <boost/compute/exception.hpp>
19 #include <boost/compute/command_queue.hpp>
20 #include <boost/compute/iterator/buffer_iterator.hpp>
21 #include <boost/compute/type_traits/type_name.hpp>
22 #include <boost/compute/detail/iterator_range_size.hpp>
23 #include <boost/compute/detail/meta_kernel.hpp>
24 #include <boost/compute/type_traits/is_device_iterator.hpp>
25
26 namespace boost {
27 namespace compute {
28 namespace detail {
29
30 template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator, class Predicate>
31 class scatter_if_kernel : meta_kernel
32 {
33 public:
scatter_if_kernel()34 scatter_if_kernel() : meta_kernel("scatter_if")
35 {}
36
set_range(InputIterator first,InputIterator last,MapIterator map,StencilIterator stencil,OutputIterator result,Predicate predicate)37 void set_range(InputIterator first,
38 InputIterator last,
39 MapIterator map,
40 StencilIterator stencil,
41 OutputIterator result,
42 Predicate predicate)
43 {
44 m_count = iterator_range_size(first, last);
45 m_input_offset = first.get_index();
46 m_output_offset = result.get_index();
47
48 m_input_offset_arg = add_arg<uint_>("input_offset");
49 m_output_offset_arg = add_arg<uint_>("output_offset");
50
51 *this <<
52 "const uint i = get_global_id(0);\n" <<
53 "uint i1 = " << map[expr<uint_>("i")] <<
54 " + output_offset;\n" <<
55 "uint i2 = i + input_offset;\n" <<
56 if_(predicate(stencil[expr<uint_>("i")])) << "\n" <<
57 result[expr<uint_>("i1")] << "=" <<
58 first[expr<uint_>("i2")] << ";\n";
59 }
60
exec(command_queue & queue)61 event exec(command_queue &queue)
62 {
63 if(m_count == 0) {
64 return event();
65 }
66
67 set_arg(m_input_offset_arg, uint_(m_input_offset));
68 set_arg(m_output_offset_arg, uint_(m_output_offset));
69
70 return exec_1d(queue, 0, m_count);
71 }
72
73 private:
74 size_t m_count;
75 size_t m_input_offset;
76 size_t m_input_offset_arg;
77 size_t m_output_offset;
78 size_t m_output_offset_arg;
79 };
80
81 } // end detail namespace
82
83 /// Copies the elements from the range [\p first, \p last) to the range
84 /// beginning at \p result using the output indices from the range beginning
85 /// at \p map if stencil is resolved to true. By default the predicate is
86 /// an identity
87 ///
88 /// Space complexity: \Omega(1)
89 template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator,
90 class Predicate>
scatter_if(InputIterator first,InputIterator last,MapIterator map,StencilIterator stencil,OutputIterator result,Predicate predicate,command_queue & queue=system::default_queue ())91 inline void scatter_if(InputIterator first,
92 InputIterator last,
93 MapIterator map,
94 StencilIterator stencil,
95 OutputIterator result,
96 Predicate predicate,
97 command_queue &queue = system::default_queue())
98 {
99 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
100 BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
101 BOOST_STATIC_ASSERT(is_device_iterator<StencilIterator>::value);
102 BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
103
104 detail::scatter_if_kernel<InputIterator, MapIterator, StencilIterator, OutputIterator, Predicate> kernel;
105
106 kernel.set_range(first, last, map, stencil, result, predicate);
107 kernel.exec(queue);
108 }
109
110 template<class InputIterator, class MapIterator, class StencilIterator, class OutputIterator>
scatter_if(InputIterator first,InputIterator last,MapIterator map,StencilIterator stencil,OutputIterator result,command_queue & queue=system::default_queue ())111 inline void scatter_if(InputIterator first,
112 InputIterator last,
113 MapIterator map,
114 StencilIterator stencil,
115 OutputIterator result,
116 command_queue &queue = system::default_queue())
117 {
118 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
119 BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
120 BOOST_STATIC_ASSERT(is_device_iterator<StencilIterator>::value);
121 BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
122
123 typedef typename std::iterator_traits<StencilIterator>::value_type T;
124
125 scatter_if(first, last, map, stencil, result, identity<T>(), queue);
126 }
127
128 } // end compute namespace
129 } // end boost namespace
130
131 #endif // BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP
132