algorithm/detail/binary_find.hpp

0001 //---------------------------------------------------------------------------//
0002 // Copyright (c) 2014 Roshan <thisisroshansmail@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_DETAIL_BINARY_FIND_HPP
0012 #define BOOST_COMPUTE_ALGORITHM_DETAIL_BINARY_FIND_HPP
0013
0014 #include <boost/compute/functional.hpp>
0015 #include <boost/compute/algorithm/find_if.hpp>
0016 #include <boost/compute/algorithm/transform.hpp>
0017 #include <boost/compute/command_queue.hpp>
0018 #include <boost/compute/detail/parameter_cache.hpp>
0019
0020 namespace boost {
0021 namespace compute {
0022 namespace detail{
0023
0024 ///
0025 /// \brief Binary find kernel class
0026 ///
0027 /// Subclass of meta_kernel to perform single step in binary find.
0028 ///
0029 template<class InputIterator, class UnaryPredicate>
0030 class binary_find_kernel : public meta_kernel
0031 {
0032 public:
0033     binary_find_kernel(InputIterator first,
0034                        InputIterator last,
0035                        UnaryPredicate predicate)
0036         : meta_kernel("binary_find")
0037     {
0038         typedef typename std::iterator_traits<InputIterator>::value_type value_type;
0039
0040         m_index_arg = add_arg<uint_ *>(memory_object::global_memory, "index");
0041         m_block_arg = add_arg<uint_>("block");
0042
0043         atomic_min<uint_> atomic_min_uint;
0044
0045         *this <<
0046             "uint i = get_global_id(0) * block;\n" <<
0047             decl<value_type>("value") << "=" << first[var<uint_>("i")] << ";\n" <<
0048             "if(" << predicate(var<value_type>("value")) << ") {\n" <<
0049                 atomic_min_uint(var<uint_ *>("index"), var<uint_>("i")) << ";\n" <<
0050             "}\n";
0051     }
0052
0053     size_t m_index_arg;
0054     size_t m_block_arg;
0055 };
0056
0057 ///
0058 /// \brief Binary find algorithm
0059 ///
0060 /// Finds the end of true values in the partitioned range [first, last).
0061 /// \return Iterator pointing to end of true values
0062 ///
0063 /// \param first Iterator pointing to start of range
0064 /// \param last Iterator pointing to end of range
0065 /// \param predicate Predicate according to which the range is partitioned
0066 /// \param queue Queue on which to execute
0067 ///
0068 template<class InputIterator, class UnaryPredicate>
0069 inline InputIterator binary_find(InputIterator first,
0070                                  InputIterator last,
0071                                  UnaryPredicate predicate,
0072                                  command_queue &queue = system::default_queue())
0073 {
0074     const device &device = queue.get_device();
0075
0076     boost::shared_ptr<parameter_cache> parameters =
0077         detail::parameter_cache::get_global_cache(device);
0078
0079     const std::string cache_key = "__boost_binary_find";
0080
0081     size_t find_if_limit = 128;
0082     size_t threads = parameters->get(cache_key, "tpb", 128);
0083     size_t count = iterator_range_size(first, last);
0084
0085     InputIterator search_first = first;
0086     InputIterator search_last = last;
0087
0088     scalar<uint_> index(queue.get_context());
0089
0090     // construct and compile binary_find kernel
0091     binary_find_kernel<InputIterator, UnaryPredicate>
0092         binary_find_kernel(search_first, search_last, predicate);
0093     ::boost::compute::kernel kernel = binary_find_kernel.compile(queue.get_context());
0094
0095     // set buffer for index
0096     kernel.set_arg(binary_find_kernel.m_index_arg, index.get_buffer());
0097
0098     while(count > find_if_limit) {
0099         index.write(static_cast<uint_>(count), queue);
0100
0101         // set block and run binary_find kernel
0102         uint_ block = static_cast<uint_>((count - 1)/(threads - 1));
0103         kernel.set_arg(binary_find_kernel.m_block_arg, block);
0104         queue.enqueue_1d_range_kernel(kernel, 0, threads, 0);
0105
0106         size_t i = index.read(queue);
0107
0108         if(i == count) {
0109             search_first = search_last - ((count - 1)%(threads - 1));
0110             break;
0111         } else {
0112             search_last = search_first + i;
0113             search_first = search_last - ((count - 1)/(threads - 1));
0114         }
0115
0116         // Make sure that first and last stay within the input range
0117         search_last = (std::min)(search_last, last);
0118         search_last = (std::max)(search_last, first);
0119
0120         search_first = (std::max)(search_first, first);
0121         search_first = (std::min)(search_first, last);
0122
0123         count = iterator_range_size(search_first, search_last);
0124     }
0125
0126     return find_if(search_first, search_last, predicate, queue);
0127 }
0128
0129 } // end detail namespace
0130 } // end compute namespace
0131 } // end boost namespace
0132
0133 #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_BINARY_FIND_HPP