1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2015 Jakub Pola <jakub.pola@gmail.com>
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
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
11 #ifndef BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP
12 #define BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP
14 #include <boost/algorithm/string/replace.hpp>
16 #include <boost/compute/system.hpp>
17 #include <boost/compute/exception.hpp>
18 #include <boost/compute/command_queue.hpp>
19 #include <boost/compute/iterator/buffer_iterator.hpp>
20 #include <boost/compute/type_traits/type_name.hpp>
21 #include <boost/compute/detail/iterator_range_size.hpp>
22 #include <boost/compute/detail/meta_kernel.hpp>
28 template<class InputIterator
, class MapIterator
, class StencilIterator
, class OutputIterator
, class Predicate
>
29 class scatter_if_kernel
: meta_kernel
32 scatter_if_kernel() : meta_kernel("scatter_if")
35 void set_range(InputIterator first
,
38 StencilIterator stencil
,
39 OutputIterator result
,
42 m_count
= iterator_range_size(first
, last
);
43 m_input_offset
= first
.get_index();
44 m_output_offset
= result
.get_index();
46 m_input_offset_arg
= add_arg
<uint_
>("input_offset");
47 m_output_offset_arg
= add_arg
<uint_
>("output_offset");
50 "const uint i = get_global_id(0);\n" <<
51 "uint i1 = " << map
[expr
<uint_
>("i")] <<
52 " + output_offset;\n" <<
53 "uint i2 = i + input_offset;\n" <<
54 if_(predicate(stencil
[expr
<uint_
>("i")])) << "\n" <<
55 result
[expr
<uint_
>("i1")] << "=" <<
56 first
[expr
<uint_
>("i2")] << ";\n";
59 event
exec(command_queue
&queue
)
65 set_arg(m_input_offset_arg
, uint_(m_input_offset
));
66 set_arg(m_output_offset_arg
, uint_(m_output_offset
));
68 return exec_1d(queue
, 0, m_count
);
73 size_t m_input_offset
;
74 size_t m_input_offset_arg
;
75 size_t m_output_offset
;
76 size_t m_output_offset_arg
;
79 } // end detail namespace
81 /// Copies the elements from the range [\p first, \p last) to the range
82 /// beginning at \p result using the output indices from the range beginning
83 /// at \p map if stencil is resolved to true. By default the predicate is
87 template<class InputIterator
, class MapIterator
, class StencilIterator
, class OutputIterator
,
89 inline void scatter_if(InputIterator first
,
92 StencilIterator stencil
,
93 OutputIterator result
,
95 command_queue
&queue
= system::default_queue())
97 detail::scatter_if_kernel
<InputIterator
, MapIterator
, StencilIterator
, OutputIterator
, Predicate
> kernel
;
99 kernel
.set_range(first
, last
, map
, stencil
, result
, predicate
);
103 template<class InputIterator
, class MapIterator
, class StencilIterator
, class OutputIterator
>
104 inline void scatter_if(InputIterator first
,
107 StencilIterator stencil
,
108 OutputIterator result
,
109 command_queue
&queue
= system::default_queue())
111 typedef typename
std::iterator_traits
<StencilIterator
>::value_type T
;
113 scatter_if(first
, last
, map
, stencil
, result
, identity
<T
>(), queue
);
116 } // end compute namespace
117 } // end boost namespace
119 #endif // BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP