File indexing completed on 2025-01-18 09:29:55
0001
0002
0003
0004
0005
0006
0007
0008
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 }
0057
0058
0059
0060
0061
0062
0063
0064
0065
0066
0067
0068
0069
0070
0071
0072
0073
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
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 }
0124 }
0125
0126 #endif