compute/algorithm/transform_if.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013-2015 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_TRANSFORM_IF_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP
0013
0014 #include <boost/static_assert.hpp>
0015
0016 #include <boost/compute/cl.hpp>
0017 #include <boost/compute/system.hpp>
0018 #include <boost/compute/command_queue.hpp>
0019 #include <boost/compute/algorithm/count.hpp>
0020 #include <boost/compute/algorithm/count_if.hpp>
0021 #include <boost/compute/algorithm/exclusive_scan.hpp>
0022 #include <boost/compute/container/vector.hpp>
0023 #include <boost/compute/detail/meta_kernel.hpp>
0024 #include <boost/compute/detail/iterator_range_size.hpp>
0025 #include <boost/compute/iterator/discard_iterator.hpp>
0026 #include <boost/compute/type_traits/is_device_iterator.hpp>
0027
0028 namespace boost {
0029 namespace compute {
0030 namespace detail {
0031
0032 // Space complexity: O(2n)
0033 template<class InputIterator, class OutputIterator, class UnaryFunction, class Predicate>
0034 inline OutputIterator transform_if_impl(InputIterator first,
0035                                         InputIterator last,
0036                                         OutputIterator result,
0037                                         UnaryFunction function,
0038                                         Predicate predicate,
0039                                         bool copyIndex,
0040                                         command_queue &queue)
0041 {
0042     typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
0043
0044     size_t count = detail::iterator_range_size(first, last);
0045     if(count == 0){
0046         return result;
0047     }
0048
0049     const context &context = queue.get_context();
0050
0051     // storage for destination indices
0052     ::boost::compute::vector<cl_uint> indices(count, context);
0053
0054     // write counts
0055     ::boost::compute::detail::meta_kernel k1("transform_if_write_counts");
0056     k1 << indices.begin()[k1.get_global_id(0)] << " = "
0057            << predicate(first[k1.get_global_id(0)]) << " ? 1 : 0;\n";
0058     k1.exec_1d(queue, 0, count);
0059
0060     // scan indices
0061     size_t copied_element_count = (indices.cend() - 1).read(queue);
0062     ::boost::compute::exclusive_scan(
0063         indices.begin(), indices.end(), indices.begin(), queue
0064     );
0065     copied_element_count += (indices.cend() - 1).read(queue); // last scan element plus last mask element
0066
0067     // copy values
0068     ::boost::compute::detail::meta_kernel k2("transform_if_do_copy");
0069     k2 << "if(" << predicate(first[k2.get_global_id(0)]) << ")" <<
0070           "    " << result[indices.begin()[k2.get_global_id(0)]] << "=";
0071
0072     if(copyIndex){
0073         k2 << k2.get_global_id(0) << ";\n";
0074     }
0075     else {
0076         k2 << function(first[k2.get_global_id(0)]) << ";\n";
0077     }
0078
0079     k2.exec_1d(queue, 0, count);
0080
0081     return result + static_cast<difference_type>(copied_element_count);
0082 }
0083
0084 template<class InputIterator, class UnaryFunction, class Predicate>
0085 inline discard_iterator transform_if_impl(InputIterator first,
0086                                           InputIterator last,
0087                                           discard_iterator result,
0088                                           UnaryFunction function,
0089                                           Predicate predicate,
0090                                           bool copyIndex,
0091                                           command_queue &queue)
0092 {
0093     (void) function;
0094     (void) copyIndex;
0095
0096     return result + count_if(first, last, predicate, queue);
0097 }
0098
0099 } // end detail namespace
0100
0101 /// Copies each element in the range [\p first, \p last) for which
0102 /// \p predicate returns \c true to the range beginning at \p result.
0103 ///
0104 /// Space complexity: O(2n)
0105 template<class InputIterator, class OutputIterator, class UnaryFunction, class Predicate>
0106 inline OutputIterator transform_if(InputIterator first,
0107                                    InputIterator last,
0108                                    OutputIterator result,
0109                                    UnaryFunction function,
0110                                    Predicate predicate,
0111                                    command_queue &queue = system::default_queue())
0112 {
0113     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0114     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
0115     return detail::transform_if_impl(
0116         first, last, result, function, predicate, false, queue
0117     );
0118 }
0119
0120 } // end compute namespace
0121 } // end boost namespace
0122
0123 #endif // BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP