Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2025-01-18 09:29:55

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2013-2014 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_ADJACENT_DIFFERENCE_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_ADJACENT_DIFFERENCE_HPP
0013 
0014 #include <iterator>
0015 
0016 #include <boost/static_assert.hpp>
0017 
0018 #include <boost/compute/system.hpp>
0019 #include <boost/compute/command_queue.hpp>
0020 #include <boost/compute/detail/meta_kernel.hpp>
0021 #include <boost/compute/detail/iterator_range_size.hpp>
0022 #include <boost/compute/functional/operator.hpp>
0023 #include <boost/compute/container/vector.hpp>
0024 #include <boost/compute/type_traits/is_device_iterator.hpp>
0025 
0026 namespace boost {
0027 namespace compute {
0028 
0029 namespace detail {
0030 
0031 template<class InputIterator, class OutputIterator, class BinaryFunction>
0032 inline OutputIterator
0033 dispatch_adjacent_difference(InputIterator first,
0034                              InputIterator last,
0035                              OutputIterator result,
0036                              BinaryFunction op,
0037                              command_queue &queue = system::default_queue())
0038 {
0039     size_t count = detail::iterator_range_size(first, last);
0040     detail::meta_kernel k("adjacent_difference");
0041 
0042     k << "const uint i = get_global_id(0);\n"
0043       << "if(i == 0){\n"
0044       << "    " << result[k.var<uint_>("0")] << " = " << first[k.var<uint_>("0")] << ";\n"
0045       << "}\n"
0046       << "else {\n"
0047       << "    " << result[k.var<uint_>("i")] << " = "
0048       <<               op(first[k.var<uint_>("i")], first[k.var<uint_>("i-1")]) << ";\n"
0049       << "}\n";
0050 
0051     k.exec_1d(queue, 0, count, 1);
0052 
0053     return result + count;
0054 }
0055 
0056 } // end detail namespace
0057 
0058 /// Stores the difference of each pair of consecutive values in the range
0059 /// [\p first, \p last) to the range beginning at \p result. If \p op is not
0060 /// provided, \c minus<T> is used.
0061 ///
0062 /// \param first first element in the input range
0063 /// \param last last element in the input range
0064 /// \param result first element in the output range
0065 /// \param op binary difference function
0066 /// \param queue command queue to perform the operation
0067 ///
0068 /// \return \c OutputIterator to the end of the result range
0069 ///
0070 /// Space complexity: \Omega(1)<br>
0071 /// Space complexity when \p result == \p first: \Omega(n)
0072 ///
0073 /// \see adjacent_find()
0074 template<class InputIterator, class OutputIterator, class BinaryFunction>
0075 inline OutputIterator
0076 adjacent_difference(InputIterator first,
0077                     InputIterator last,
0078                     OutputIterator result,
0079                     BinaryFunction op,
0080                     command_queue &queue = system::default_queue())
0081 {
0082     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0083     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
0084     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0085 
0086     if(first == last) {
0087         return result;
0088     }
0089 
0090     if (first == result) {
0091         vector<value_type> temp(detail::iterator_range_size(first, last),
0092                                 queue.get_context());
0093         copy(first, last, temp.begin(), queue);
0094 
0095         return ::boost::compute::detail::dispatch_adjacent_difference(
0096             temp.begin(), temp.end(), result, op, queue
0097         );
0098     }
0099     else {
0100         return ::boost::compute::detail::dispatch_adjacent_difference(
0101             first, last, result, op, queue
0102         );
0103     }
0104 }
0105 
0106 /// \overload
0107 template<class InputIterator, class OutputIterator>
0108 inline OutputIterator
0109 adjacent_difference(InputIterator first,
0110                     InputIterator last,
0111                     OutputIterator result,
0112                     command_queue &queue = system::default_queue())
0113 {
0114     BOOST_STATIC_ASSERT(is_device_iterator<InputIterator>::value);
0115     BOOST_STATIC_ASSERT(is_device_iterator<OutputIterator>::value);
0116     typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0117 
0118     return ::boost::compute::adjacent_difference(
0119         first, last, result, ::boost::compute::minus<value_type>(), queue
0120     );
0121 }
0122 
0123 } // end compute namespace
0124 } // end boost namespace
0125 
0126 #endif // BOOST_COMPUTE_ALGORITHM_ADJACENT_DIFFERENCE_HPP