1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@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_GATHER_HPP
12 #define BOOST_COMPUTE_ALGORITHM_GATHER_HPP
13 
14 #include <boost/static_assert.hpp>
15 
16 #include <boost/compute/command_queue.hpp>
17 #include <boost/compute/detail/iterator_range_size.hpp>
18 #include <boost/compute/detail/meta_kernel.hpp>
19 #include <boost/compute/exception.hpp>
20 #include <boost/compute/iterator/buffer_iterator.hpp>
21 #include <boost/compute/system.hpp>
22 #include <boost/compute/type_traits/type_name.hpp>
23 #include <boost/compute/type_traits/is_device_iterator.hpp>
24 
25 namespace boost {
26 namespace compute {
27 namespace detail {
28 
29 template<class InputIterator, class MapIterator, class OutputIterator>
30 class gather_kernel : public meta_kernel
31 {
32 public:
33     gather_kernel() : meta_kernel("gather")
34     {}
35 
36     void set_range(MapIterator first,
37                    MapIterator last,
38                    InputIterator input,
39                    OutputIterator result)
40     {
41         m_count = iterator_range_size(first, last);
42 
43         *this <<
44             "const uint i = get_global_id(0);\n" <<
45             result[expr<uint_>("i")] << "=" <<
46                 input[first[expr<uint_>("i")]] << ";\n";
47     }
48 
49     event exec(command_queue &queue)
50     {
51         if(m_count == 0) {
52             return event();
53         }
54 
55         return exec_1d(queue, 0, m_count);
56     }
57 
58 private:
59     size_t m_count;
60 };
61 
62 } // end detail namespace
63 
64 /// Copies the elements using the indices from the range [\p first, \p last)
65 /// to the range beginning at \p result using the input values from the range
66 /// beginning at \p input.
67 ///
68 /// Space complexity: \Omega(1)
69 ///
70 /// \see scatter()
71 template<class InputIterator, class MapIterator, class OutputIterator>
72 inline void gather(MapIterator first,
73                    MapIterator last,
74                    InputIterator input,
75                    OutputIterator result,
76                    command_queue &queue = system::default_queue())
77 {
78     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
79     BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
80     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
81 
82     detail::gather_kernel<InputIterator, MapIterator, OutputIterator> kernel;
83 
84     kernel.set_range(first, last, input, result);
85     kernel.exec(queue);
86 }
87 
88 } // end compute namespace
89 } // end boost namespace
90 
91 #endif // BOOST_COMPUTE_ALGORITHM_GATHER_HPP
92