Back to home page

EIC code displayed by LXR

 
 

    


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

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2014 Mageswaran.D <mageswaran1989@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 #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         //Individual elements are compared and results are stored in parallel.
0036         //0 is true
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     // load (or create) lexicographical compare program
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 } // end detail namespace
0106 
0107 /// Checks if the first range [first1, last1) is lexicographically
0108 /// less than the second range [first2, last2).
0109 ///
0110 /// Space complexity:
0111 /// \Omega(max(distance(\p first1, \p last1), distance(\p first2, \p last2)))
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 } // end compute namespace
0126 } // end boost namespac