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/reduce_on_cpu.hpp (3949B)

//---------------------------------------------------------------------------//
// Copyright (c) 2016 Jakub Szuppe <j.szuppe@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_REDUCE_ON_CPU_HPP
#define BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP

#include <algorithm>

#include <boost/compute/buffer.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
#include <boost/compute/iterator/buffer_iterator.hpp>
#include <boost/compute/type_traits/result_of.hpp>
#include <boost/compute/algorithm/detail/serial_reduce.hpp>

namespace boost {
namespace compute {
namespace detail {

template<class InputIterator, class OutputIterator, class BinaryFunction>
inline void reduce_on_cpu(InputIterator first,
                          InputIterator last,
                          OutputIterator result,
                          BinaryFunction function,
                          command_queue &queue)
{
    typedef typename
        std::iterator_traits<InputIterator>::value_type T;
    typedef typename
        ::boost::compute::result_of<BinaryFunction(T, T)>::type result_type;

const device &device = queue.get_device();
    const uint_ compute_units = queue.get_device().compute_units();

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

std::string cache_key =
        "__boost_reduce_cpu_" + boost::lexical_cast<std::string>(sizeof(T));

// for inputs smaller than serial_reduce_threshold
    // serial_reduce algorithm is used
    uint_ serial_reduce_threshold =
        parameters->get(cache_key, "serial_reduce_threshold", 16384 * sizeof(T));
    serial_reduce_threshold =
        (std::max)(serial_reduce_threshold, uint_(compute_units));

const context &context = queue.get_context();
    size_t count = detail::iterator_range_size(first, last);
    if(count == 0){
        return;
    }
    else if(count < serial_reduce_threshold) {
        return serial_reduce(first, last, result, function, queue);
    }

meta_kernel k("reduce_on_cpu");
    buffer output(context, sizeof(result_type) * compute_units);

size_t count_arg = k.add_arg<uint_>("count");
    size_t output_arg =
        k.add_arg<result_type *>(memory_object::global_memory, "output");

k <<
        "uint block = " <<
            "(uint)ceil(((float)count)/get_global_size(0));\n" <<
        "uint index = get_global_id(0) * block;\n" <<
        "uint end = min(count, index + block);\n" <<

k.decl<result_type>("result") << " = " << first[k.var<uint_>("index")] << ";\n" <<
        "index++;\n" <<
        "while(index < end){\n" <<
             "result = " << function(k.var<T>("result"),
                                     first[k.var<uint_>("index")]) << ";\n" <<
             "index++;\n" <<
        "}\n" <<
        "output[get_global_id(0)] = result;\n";

size_t global_work_size = compute_units;
    kernel kernel = k.compile(context);

// reduction to global_work_size elements
    kernel.set_arg(count_arg, static_cast<uint_>(count));
    kernel.set_arg(output_arg, output);
    queue.enqueue_1d_range_kernel(kernel, 0, global_work_size, 0);

// final reduction
    reduce_on_cpu(
        make_buffer_iterator<result_type>(output),
        make_buffer_iterator<result_type>(output, global_work_size),
        result,
        function,
        queue
    );
}

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

#endif // BOOST_COMPUTE_ALGORITHM_DETAIL_REDUCE_ON_CPU_HPP

cmd: