File indexing completed on 2025-01-18 09:29:56
0001
0002
0003
0004
0005
0006
0007
0008
0009
0010
0011 #include <boost/static_assert.hpp>
0012
0013 #include <boost/compute/system.hpp>
0014 #include <boost/compute/context.hpp>
0015 #include <boost/compute/command_queue.hpp>
0016 #include <boost/compute/algorithm/any_of.hpp>
0017 #include <boost/compute/container/vector.hpp>
0018 #include <boost/compute/utility/program_cache.hpp>
0019 #include <boost/compute/type_traits/is_device_iterator.hpp>
0020
0021 namespace boost {
0022 namespace compute {
0023
0024 namespace detail {
0025
0026 const char lexicographical_compare_source[] =
0027 "__kernel void lexicographical_compare(const uint size1,\n"
0028 " const uint size2,\n"
0029 " __global const T1 *range1,\n"
0030 " __global const T2 *range2,\n"
0031 " __global bool *result_buf)\n"
0032 "{\n"
0033 " const uint i = get_global_id(0);\n"
0034 " if((i != size1) && (i != size2)){\n"
0035
0036
0037 " if(range1[i] < range2[i])\n"
0038 " result_buf[i] = 0;\n"
0039 " else\n"
0040 " result_buf[i] = 1;\n"
0041 " }\n"
0042 " else\n"
0043 " result_buf[i] = !((i == size1) && (i != size2));\n"
0044 "}\n";
0045
0046 template<class InputIterator1, class InputIterator2>
0047 inline bool dispatch_lexicographical_compare(InputIterator1 first1,
0048 InputIterator1 last1,
0049 InputIterator2 first2,
0050 InputIterator2 last2,
0051 command_queue &queue)
0052 {
0053 const boost::compute::context &context = queue.get_context();
0054
0055 boost::shared_ptr<program_cache> cache =
0056 program_cache::get_global_cache(context);
0057
0058 size_t iterator_size1 = iterator_range_size(first1, last1);
0059 size_t iterator_size2 = iterator_range_size(first2, last2);
0060 size_t max_size = (std::max)(iterator_size1, iterator_size2);
0061
0062 if(max_size == 0){
0063 return false;
0064 }
0065
0066 boost::compute::vector<bool> result_vector(max_size, context);
0067
0068
0069 typedef typename std::iterator_traits<InputIterator1>::value_type value_type1;
0070 typedef typename std::iterator_traits<InputIterator2>::value_type value_type2;
0071
0072
0073 std::string cache_key =
0074 std::string("__boost_lexicographical_compare")
0075 + type_name<value_type1>() + type_name<value_type2>();
0076
0077 std::stringstream options;
0078 options << " -DT1=" << type_name<value_type1>();
0079 options << " -DT2=" << type_name<value_type2>();
0080
0081 program lexicographical_compare_program = cache->get_or_build(
0082 cache_key, options.str(), lexicographical_compare_source, context
0083 );
0084
0085 kernel lexicographical_compare_kernel(lexicographical_compare_program,
0086 "lexicographical_compare");
0087
0088 lexicographical_compare_kernel.set_arg<uint_>(0, iterator_size1);
0089 lexicographical_compare_kernel.set_arg<uint_>(1, iterator_size2);
0090 lexicographical_compare_kernel.set_arg(2, first1.get_buffer());
0091 lexicographical_compare_kernel.set_arg(3, first2.get_buffer());
0092 lexicographical_compare_kernel.set_arg(4, result_vector.get_buffer());
0093
0094 queue.enqueue_1d_range_kernel(lexicographical_compare_kernel,
0095 0,
0096 max_size,
0097 0);
0098
0099 return boost::compute::any_of(result_vector.begin(),
0100 result_vector.end(),
0101 _1 == 0,
0102 queue);
0103 }
0104
0105 }
0106
0107
0108
0109
0110
0111
0112 template<class InputIterator1, class InputIterator2>
0113 inline bool lexicographical_compare(InputIterator1 first1,
0114 InputIterator1 last1,
0115 InputIterator2 first2,
0116 InputIterator2 last2,
0117 command_queue &queue = system::default_queue())
0118 {
0119 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator1>::value);
0120 BOOST_STATIC_ASSERT(is_device_iterator<InputIterator2>::value);
0121
0122 return detail::dispatch_lexicographical_compare(first1, last1, first2, last2, queue);
0123 }
0124
0125 }
0126 }