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