//---------------------------------------------------------------------------// // Copyright (c) 2013 Kyle Lutz // // Distributed under the Boost Software License, Version 1.0 // See accompanying file LICENSE_1_0.txt or copy at // http://www.boost.org/LICENSE_1_0.txt // // See http://boostorg.github.com/compute for more information. //---------------------------------------------------------------------------// #ifndef BOOST_COMPUTE_ALGORITHM_SCATTER_HPP #define BOOST_COMPUTE_ALGORITHM_SCATTER_HPP #include #include #include #include #include #include #include #include namespace boost { namespace compute { namespace detail { template class scatter_kernel : meta_kernel { public: scatter_kernel() : meta_kernel("scatter") {} void set_range(InputIterator first, InputIterator last, MapIterator map, OutputIterator result) { m_count = iterator_range_size(first, last); m_input_offset = first.get_index(); m_output_offset = result.get_index(); m_input_offset_arg = add_arg("input_offset"); m_output_offset_arg = add_arg("output_offset"); *this << "const uint i = get_global_id(0);\n" << "uint i1 = " << map[expr("i")] << " + output_offset;\n" << "uint i2 = i + input_offset;\n" << result[expr("i1")] << "=" << first[expr("i2")] << ";\n"; } event exec(command_queue &queue) { if(m_count == 0) { return event(); } set_arg(m_input_offset_arg, uint_(m_input_offset)); set_arg(m_output_offset_arg, uint_(m_output_offset)); return exec_1d(queue, 0, m_count); } private: size_t m_count; size_t m_input_offset; size_t m_input_offset_arg; size_t m_output_offset; size_t m_output_offset_arg; }; } // end detail namespace /// Copies the elements from the range [\p first, \p last) to the range /// beginning at \p result using the output indices from the range beginning /// at \p map. /// /// \see gather() template inline void scatter(InputIterator first, InputIterator last, MapIterator map, OutputIterator result, command_queue &queue = system::default_queue()) { detail::scatter_kernel kernel; kernel.set_range(first, last, map, result); kernel.exec(queue); } } // end compute namespace } // end boost namespace #endif // BOOST_COMPUTE_ALGORITHM_SCATTER_HPP