//---------------------------------------------------------------------------// // Copyright (c) 2015 Jakub Pola // // 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_IF_HPP #define BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP #include #include #include #include #include #include #include #include #include #include namespace boost { namespace compute { namespace detail { template class scatter_if_kernel : meta_kernel { public: scatter_if_kernel() : meta_kernel("scatter_if") {} void set_range(InputIterator first, InputIterator last, MapIterator map, StencilIterator stencil, OutputIterator result, Predicate predicate) { 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" << if_(predicate(stencil[expr("i")])) << "\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 if stencil is resolved to true. By default the predicate is /// an identity /// /// Space complexity: \Omega(1) template inline void scatter_if(InputIterator first, InputIterator last, MapIterator map, StencilIterator stencil, OutputIterator result, Predicate predicate, command_queue &queue = system::default_queue()) { BOOST_STATIC_ASSERT(is_device_iterator::value); BOOST_STATIC_ASSERT(is_device_iterator::value); BOOST_STATIC_ASSERT(is_device_iterator::value); BOOST_STATIC_ASSERT(is_device_iterator::value); detail::scatter_if_kernel kernel; kernel.set_range(first, last, map, stencil, result, predicate); kernel.exec(queue); } template inline void scatter_if(InputIterator first, InputIterator last, MapIterator map, StencilIterator stencil, OutputIterator result, command_queue &queue = system::default_queue()) { BOOST_STATIC_ASSERT(is_device_iterator::value); BOOST_STATIC_ASSERT(is_device_iterator::value); BOOST_STATIC_ASSERT(is_device_iterator::value); BOOST_STATIC_ASSERT(is_device_iterator::value); typedef typename std::iterator_traits::value_type T; scatter_if(first, last, map, stencil, result, identity(), queue); } } // end compute namespace } // end boost namespace #endif // BOOST_COMPUTE_ALGORITHM_SCATTER_IF_HPP