File indexing completed on 2025-01-18 09:29:53
0001
0002
0003
0004
0005
0006
0007
0008
0009
0010
0011 #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_BALANCED_PATH_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_DETAIL_BALANCED_PATH_HPP
0013
0014 #include <iterator>
0015
0016 #include <boost/compute/algorithm/find_if.hpp>
0017 #include <boost/compute/container/vector.hpp>
0018 #include <boost/compute/detail/iterator_range_size.hpp>
0019 #include <boost/compute/detail/meta_kernel.hpp>
0020 #include <boost/compute/lambda.hpp>
0021 #include <boost/compute/system.hpp>
0022
0023 namespace boost {
0024 namespace compute {
0025 namespace detail {
0026
0027
0028
0029
0030
0031
0032
0033 class balanced_path_kernel : public meta_kernel
0034 {
0035 public:
0036 unsigned int tile_size;
0037
0038 balanced_path_kernel() : meta_kernel("balanced_path")
0039 {
0040 tile_size = 4;
0041 }
0042
0043 template<class InputIterator1, class InputIterator2,
0044 class OutputIterator1, class OutputIterator2,
0045 class Compare>
0046 void set_range(InputIterator1 first1,
0047 InputIterator1 last1,
0048 InputIterator2 first2,
0049 InputIterator2 last2,
0050 OutputIterator1 result_a,
0051 OutputIterator2 result_b,
0052 Compare comp)
0053 {
0054 typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
0055
0056 m_a_count = iterator_range_size(first1, last1);
0057 m_a_count_arg = add_arg<uint_>("a_count");
0058
0059 m_b_count = iterator_range_size(first2, last2);
0060 m_b_count_arg = add_arg<uint_>("b_count");
0061
0062 *this <<
0063 "uint i = get_global_id(0);\n" <<
0064 "uint target = (i+1)*" << tile_size << ";\n" <<
0065 "uint start = max(convert_int(0),convert_int(target)-convert_int(b_count));\n" <<
0066 "uint end = min(target,a_count);\n" <<
0067 "uint a_index, b_index;\n" <<
0068 "while(start<end)\n" <<
0069 "{\n" <<
0070 " a_index = (start + end)/2;\n" <<
0071 " b_index = target - a_index - 1;\n" <<
0072 " if(!(" << comp(first2[expr<uint_>("b_index")],
0073 first1[expr<uint_>("a_index")]) << "))\n" <<
0074 " start = a_index + 1;\n" <<
0075 " else end = a_index;\n" <<
0076 "}\n" <<
0077 "a_index = start;\n" <<
0078 "b_index = target - start;\n" <<
0079 "if(b_index < b_count)\n" <<
0080 "{\n" <<
0081 " " << decl<const value_type>("x") << " = " <<
0082 first2[expr<uint_>("b_index")] << ";\n" <<
0083 " uint a_start = 0, a_end = a_index, a_mid;\n" <<
0084 " uint b_start = 0, b_end = b_index, b_mid;\n" <<
0085 " while(a_start<a_end)\n" <<
0086 " {\n" <<
0087 " a_mid = (a_start + a_end)/2;\n" <<
0088 " if(" << comp(first1[expr<uint_>("a_mid")], expr<value_type>("x")) << ")\n" <<
0089 " a_start = a_mid+1;\n" <<
0090 " else a_end = a_mid;\n" <<
0091 " }\n" <<
0092 " while(b_start<b_end)\n" <<
0093 " {\n" <<
0094 " b_mid = (b_start + b_end)/2;\n" <<
0095 " if(" << comp(first2[expr<uint_>("b_mid")], expr<value_type>("x")) << ")\n" <<
0096 " b_start = b_mid+1;\n" <<
0097 " else b_end = b_mid;\n" <<
0098 " }\n" <<
0099 " uint a_run = a_index - a_start;\n" <<
0100 " uint b_run = b_index - b_start;\n" <<
0101 " uint x_count = a_run + b_run;\n" <<
0102 " uint b_advance = max(x_count / 2, x_count - a_run);\n" <<
0103 " b_end = min(b_count, b_start + b_advance + 1);\n" <<
0104 " uint temp_start = b_index, temp_end = b_end, temp_mid;" <<
0105 " while(temp_start < temp_end)\n" <<
0106 " {\n" <<
0107 " temp_mid = (temp_start + temp_end + 1)/2;\n" <<
0108 " if(" << comp(expr<value_type>("x"), first2[expr<uint_>("temp_mid")]) << ")\n" <<
0109 " temp_end = temp_mid-1;\n" <<
0110 " else temp_start = temp_mid;\n" <<
0111 " }\n" <<
0112 " b_run = temp_start - b_start + 1;\n" <<
0113 " b_advance = min(b_advance, b_run);\n" <<
0114 " uint a_advance = x_count - b_advance;\n" <<
0115 " uint star = convert_uint((a_advance == b_advance + 1) " <<
0116 "&& (b_advance < b_run));\n" <<
0117 " a_index = a_start + a_advance;\n" <<
0118 " b_index = target - a_index + star;\n" <<
0119 "}\n" <<
0120 result_a[expr<uint_>("i")] << " = a_index;\n" <<
0121 result_b[expr<uint_>("i")] << " = b_index;\n";
0122
0123 }
0124
0125 template<class InputIterator1, class InputIterator2,
0126 class OutputIterator1, class OutputIterator2>
0127 void set_range(InputIterator1 first1,
0128 InputIterator1 last1,
0129 InputIterator2 first2,
0130 InputIterator2 last2,
0131 OutputIterator1 result_a,
0132 OutputIterator2 result_b)
0133 {
0134 typedef typename std::iterator_traits<InputIterator1>::value_type value_type;
0135 ::boost::compute::less<value_type> less_than;
0136 set_range(first1, last1, first2, last2, result_a, result_b, less_than);
0137 }
0138
0139 event exec(command_queue &queue)
0140 {
0141 if((m_a_count + m_b_count)/tile_size == 0) {
0142 return event();
0143 }
0144
0145 set_arg(m_a_count_arg, uint_(m_a_count));
0146 set_arg(m_b_count_arg, uint_(m_b_count));
0147
0148 return exec_1d(queue, 0, (m_a_count + m_b_count)/tile_size);
0149 }
0150
0151 private:
0152 size_t m_a_count;
0153 size_t m_a_count_arg;
0154 size_t m_b_count;
0155 size_t m_b_count_arg;
0156 };
0157
0158 }
0159 }
0160 }
0161
0162 #endif