compute/algorithm/gather.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
0003 //
0004 // Distributed under the Boost Software License, Version 1.0
0005 // See accompanying file LICENSE_1_0.txt or copy at
0006 // http://www.boost.org/LICENSE_1_0.txt
0007 //
0008 // See http://boostorg.github.com/compute for more information.
0009 //---------------------------------------------------------------------------//
0010
0011 #ifndef BOOST_COMPUTE_ALGORITHM_GATHER_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_GATHER_HPP
0013
0014 #include <boost/static_assert.hpp>
0015
0016 #include <boost/compute/command_queue.hpp>
0017 #include <boost/compute/detail/iterator_range_size.hpp>
0018 #include <boost/compute/detail/meta_kernel.hpp>
0019 #include <boost/compute/exception.hpp>
0020 #include <boost/compute/iterator/buffer_iterator.hpp>
0021 #include <boost/compute/system.hpp>
0022 #include <boost/compute/type_traits/type_name.hpp>
0023 #include <boost/compute/type_traits/is_device_iterator.hpp>
0024
0025 namespace boost {
0026 namespace compute {
0027 namespace detail {
0028
0029 template<class InputIterator, class MapIterator, class OutputIterator>
0030 class gather_kernel : public meta_kernel
0031 {
0032 public:
0033     gather_kernel() : meta_kernel("gather")
0034     {}
0035
0036     void set_range(MapIterator first,
0037                    MapIterator last,
0038                    InputIterator input,
0039                    OutputIterator result)
0040     {
0041         m_count = iterator_range_size(first, last);
0042
0043         *this <<
0044             "const uint i = get_global_id(0);\n" <<
0045             result[expr<uint_>("i")] << "=" <<
0046                 input[first[expr<uint_>("i")]] << ";\n";
0047     }
0048
0049     event exec(command_queue &queue)
0050     {
0051         if(m_count == 0) {
0052             return event();
0053         }
0054
0055         return exec_1d(queue, 0, m_count);
0056     }
0057
0058 private:
0059     size_t m_count;
0060 };
0061
0062 } // end detail namespace
0063
0064 /// Copies the elements using the indices from the range [\p first, \p last)
0065 /// to the range beginning at \p result using the input values from the range
0066 /// beginning at \p input.
0067 ///
0068 /// Space complexity: \Omega(1)
0069 ///
0070 /// \see scatter()
0071 template<class InputIterator, class MapIterator, class OutputIterator>
0072 inline void gather(MapIterator first,
0073                    MapIterator last,
0074                    InputIterator input,
0075                    OutputIterator result,
0076                    command_queue &queue = system::default_queue())
0077 {
0078     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0079     BOOST_STATIC_ASSERT(is_device_iterator<MapIterator>::value);
0080     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
0081
0082     detail::gather_kernel<InputIterator, MapIterator, OutputIterator> kernel;
0083
0084     kernel.set_range(first, last, input, result);
0085     kernel.exec(queue);
0086 }
0087
0088 } // end compute namespace
0089 } // end boost namespace
0090
0091 #endif // BOOST_COMPUTE_ALGORITHM_GATHER_HPP