FM / detail

/usr/include/boost/compute/algorithm/detail/

/usr/include/boost/compute/algorithm/detail

Name	Size	Mode	Actions
balanced_path.hpp	6267	0644	edit dl rm
binary_find.hpp	4626	0644	edit dl rm
compact.hpp	2188	0644	edit dl rm
copy_on_device.hpp	7382	0644	edit dl rm
copy_to_device.hpp	6731	0644	edit dl rm
copy_to_host.hpp	6658	0644	edit dl rm
count_if_with_ballot.hpp	2605	0644	edit dl rm
count_if_with_reduce.hpp	2648	0644	edit dl rm
count_if_with_threads.hpp	4241	0644	edit dl rm
find_extrema.hpp	2441	0644	edit dl rm
find_extrema_on_cpu.hpp	5333	0644	edit dl rm
find_extrema_with_atomics.hpp	4251	0644	edit dl rm
find_extrema_with_reduce.hpp	18672	0644	edit dl rm
find_if_with_atomics.hpp	8509	0644	edit dl rm
inplace_reduce.hpp	4839	0644	edit dl rm
insertion_sort.hpp	6030	0644	edit dl rm
merge_path.hpp	3874	0644	edit dl rm
merge_sort_on_cpu.hpp	14654	0644	edit dl rm
merge_sort_on_gpu.hpp	22093	0644	edit dl rm
merge_with_merge_path.hpp	7339	0644	edit dl rm
radix_sort.hpp	15502	0644	edit dl rm
random_fill.hpp	1880	0644	edit dl rm
reduce_by_key.hpp	5243	0644	edit dl rm
reduce_by_key_with_scan.hpp	23735	0644	edit dl rm
reduce_on_cpu.hpp	3949	0644	edit dl rm
reduce_on_gpu.hpp	10391	0644	edit dl rm
scan.hpp	1545	0644	edit dl rm
scan_on_cpu.hpp	7056	0644	edit dl rm
scan_on_gpu.hpp	11320	0644	edit dl rm
search_all.hpp	2579	0644	edit dl rm
serial_accumulate.hpp	1990	0644	edit dl rm
serial_count_if.hpp	2243	0644	edit dl rm
serial_find_extrema.hpp	3151	0644	edit dl rm
serial_merge.hpp	3504	0644	edit dl rm
serial_reduce.hpp	2145	0644	edit dl rm
serial_reduce_by_key.hpp	4240	0644	edit dl rm
serial_scan.hpp	3184	0644	edit dl rm

Edit: /usr/include/boost/compute/algorithm/detail/binary_find.hpp (4626B)

//---------------------------------------------------------------------------//
// Copyright (c) 2014 Roshan <thisisroshansmail@gmail.com>
//
// Distributed under the Boost Software License, Version 1.0
// See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt
//
// See http://boostorg.github.com/compute for more information.
//---------------------------------------------------------------------------//

#ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_BINARY_FIND_HPP
#define BOOST_COMPUTE_ALGORITHM_DETAIL_BINARY_FIND_HPP

#include <boost/compute/functional.hpp>
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/algorithm/transform.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/parameter_cache.hpp>

namespace boost {
namespace compute {
namespace detail{

///
/// \brief Binary find kernel class
///
/// Subclass of meta_kernel to perform single step in binary find.
///
template<class InputIterator, class UnaryPredicate>
class binary_find_kernel : public meta_kernel
{
public:
    binary_find_kernel(InputIterator first,
                       InputIterator last,
                       UnaryPredicate predicate)
        : meta_kernel("binary_find")
    {
        typedef typename std::iterator_traits<InputIterator>::value_type value_type;

m_index_arg = add_arg<uint_ *>(memory_object::global_memory, "index");
        m_block_arg = add_arg<uint_>("block");

atomic_min<uint_> atomic_min_uint;

*this <<
            "uint i = get_global_id(0) * block;\n" <<
            decl<value_type>("value") << "=" << first[var<uint_>("i")] << ";\n" <<
            "if(" << predicate(var<value_type>("value")) << ") {\n" <<
                atomic_min_uint(var<uint_ *>("index"), var<uint_>("i")) << ";\n" <<
            "}\n";
    }

size_t m_index_arg;
    size_t m_block_arg;
};

///
/// \brief Binary find algorithm
///
/// Finds the end of true values in the partitioned range [first, last).
/// \return Iterator pointing to end of true values
///
/// \param first Iterator pointing to start of range
/// \param last Iterator pointing to end of range
/// \param predicate Predicate according to which the range is partitioned
/// \param queue Queue on which to execute
///
template<class InputIterator, class UnaryPredicate>
inline InputIterator binary_find(InputIterator first,
                                 InputIterator last,
                                 UnaryPredicate predicate,
                                 command_queue &queue = system::default_queue())
{
    const device &device = queue.get_device();

boost::shared_ptr<parameter_cache> parameters =
        detail::parameter_cache::get_global_cache(device);

const std::string cache_key = "__boost_binary_find";

size_t find_if_limit = 128;
    size_t threads = parameters->get(cache_key, "tpb", 128);
    size_t count = iterator_range_size(first, last);

InputIterator search_first = first;
    InputIterator search_last = last;

scalar<uint_> index(queue.get_context());

// construct and compile binary_find kernel
    binary_find_kernel<InputIterator, UnaryPredicate>
        binary_find_kernel(search_first, search_last, predicate);
    ::boost::compute::kernel kernel = binary_find_kernel.compile(queue.get_context());

// set buffer for index
    kernel.set_arg(binary_find_kernel.m_index_arg, index.get_buffer());

while(count > find_if_limit) {
        index.write(static_cast<uint_>(count), queue);

// set block and run binary_find kernel
        uint_ block = static_cast<uint_>((count - 1)/(threads - 1));
        kernel.set_arg(binary_find_kernel.m_block_arg, block);
        queue.enqueue_1d_range_kernel(kernel, 0, threads, 0);

size_t i = index.read(queue);

if(i == count) {
            search_first = search_last - ((count - 1)%(threads - 1));
            break;
        } else {
            search_last = search_first + i;
            search_first = search_last - ((count - 1)/(threads - 1));
        }

// Make sure that first and last stay within the input range
        search_last = (std::min)(search_last, last);
        search_last = (std::max)(search_last, first);

search_first = (std::max)(search_first, first);
        search_first = (std::min)(search_first, last);

count = iterator_range_size(search_first, search_last);
    }

return find_if(search_first, search_last, predicate, queue);
}

} // end detail namespace
} // end compute namespace
} // end boost namespace

#endif // BOOST_COMPUTE_ALGORITHM_DETAIL_BINARY_FIND_HPP

cmd: