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/copy_on_device.hpp (7382B)

//---------------------------------------------------------------------------//
// Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@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_COPY_ON_DEVICE_HPP
#define BOOST_COMPUTE_ALGORITHM_DETAIL_COPY_ON_DEVICE_HPP

#include <iterator>

#include <boost/compute/command_queue.hpp>
#include <boost/compute/async/future.hpp>
#include <boost/compute/iterator/buffer_iterator.hpp>
#include <boost/compute/iterator/discard_iterator.hpp>
#include <boost/compute/memory/svm_ptr.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
#include <boost/compute/detail/work_size.hpp>
#include <boost/compute/detail/vendor.hpp>

namespace boost {
namespace compute {
namespace detail {

template<class InputIterator, class OutputIterator>
inline event copy_on_device_cpu(InputIterator first,
                                OutputIterator result,
                                size_t count,
                                command_queue &queue,
                                const wait_list &events)
{
    meta_kernel k("copy");
    const device& device = queue.get_device();

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" <<
        "while(index < end){\n" <<
            result[k.var<uint_>("index")] << '=' <<
                first[k.var<uint_>("index")] << ";\n" <<
            "index++;\n" <<
        "}\n";

k.add_set_arg<const uint_>("count", static_cast<uint_>(count));

size_t global_work_size = device.compute_units();
    if(count <= 1024) global_work_size = 1;
    return k.exec_1d(queue, 0, global_work_size, events);
}

template<class InputIterator, class OutputIterator>
inline event copy_on_device_gpu(InputIterator first,
                                OutputIterator result,
                                size_t count,
                                command_queue &queue,
                                const wait_list &events)
{
    typedef typename std::iterator_traits<InputIterator>::value_type input_type;

const device& device = queue.get_device();
    boost::shared_ptr<parameter_cache> parameters =
        detail::parameter_cache::get_global_cache(device);
    std::string cache_key =
        "__boost_copy_kernel_" + boost::lexical_cast<std::string>(sizeof(input_type));

uint_ vpt = parameters->get(cache_key, "vpt", 4);
    uint_ tpb = parameters->get(cache_key, "tpb", 128);

meta_kernel k("copy");
    k <<
        "uint index = get_local_id(0) + " <<
            "(" << vpt * tpb << " * get_group_id(0));\n" <<
        "for(uint i = 0; i < " << vpt << "; i++){\n" <<
        "    if(index < count){\n" <<
                result[k.var<uint_>("index")] << '=' <<
                    first[k.var<uint_>("index")] << ";\n" <<
        "       index += " << tpb << ";\n"
        "    }\n"
        "}\n";

k.add_set_arg<const uint_>("count", static_cast<uint_>(count));
    size_t global_work_size = calculate_work_size(count, vpt, tpb);
    return k.exec_1d(queue, 0, global_work_size, tpb, events);
}

template<class InputIterator, class OutputIterator>
inline event dispatch_copy_on_device(InputIterator first,
                                     InputIterator last,
                                     OutputIterator result,
                                     command_queue &queue,
                                     const wait_list &events)
{
    const size_t count = detail::iterator_range_size(first, last);

if(count == 0){
        // nothing to do
        return event();
    }

const device& device = queue.get_device();
    // copy_on_device_cpu() does not work for CPU on Apple platform
    // due to bug in its compiler.
    // See https://github.com/boostorg/compute/pull/626
    if((device.type() & device::cpu) && !is_apple_platform_device(device))
    {
        return copy_on_device_cpu(first, result, count, queue, events);
    }
    return copy_on_device_gpu(first, result, count, queue, events);
}

template<class InputIterator, class OutputIterator>
inline OutputIterator copy_on_device(InputIterator first,
                                     InputIterator last,
                                     OutputIterator result,
                                     command_queue &queue,
                                     const wait_list &events)
{
    dispatch_copy_on_device(first, last, result, queue, events);
    return result + std::distance(first, last);
}

template<class InputIterator>
inline discard_iterator copy_on_device(InputIterator first,
                                       InputIterator last,
                                       discard_iterator result,
                                       command_queue &queue,
                                       const wait_list &events)
{
    (void) queue;
    (void) events;

return result + std::distance(first, last);
}

template<class InputIterator, class OutputIterator>
inline future<OutputIterator> copy_on_device_async(InputIterator first,
                                                   InputIterator last,
                                                   OutputIterator result,
                                                   command_queue &queue,
                                                   const wait_list &events)
{
    event event_ = dispatch_copy_on_device(first, last, result, queue, events);
    return make_future(result + std::distance(first, last), event_);
}

#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_on_device() specialization for svm_ptr
template<class T>
inline svm_ptr<T> copy_on_device(svm_ptr<T> first,
                                 svm_ptr<T> last,
                                 svm_ptr<T> result,
                                 command_queue &queue,
                                 const wait_list &events)
{
    size_t count = iterator_range_size(first, last);
    if(count == 0){
        return result;
    }

queue.enqueue_svm_memcpy(
        result.get(), first.get(), count * sizeof(T), events
    );

return result + count;
}

template<class T>
inline future<svm_ptr<T> > copy_on_device_async(svm_ptr<T> first,
                                                svm_ptr<T> last,
                                                svm_ptr<T> result,
                                                command_queue &queue,
                                                const wait_list &events)
{
    size_t count = iterator_range_size(first, last);
    if(count == 0){
        return future<svm_ptr<T> >();
    }

event event_ = queue.enqueue_svm_memcpy_async(
        result.get(), first.get(), count * sizeof(T), events
    );

return make_future(result + count, event_);
}
#endif // BOOST_COMPUTE_CL_VERSION_2_0

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

#endif // BOOST_COMPUTE_ALGORITHM_DETAIL_COPY_ON_DEVICE_HPP

//---------------------------------------------------------------------------// // Copyright (c) 2013 Kyle Lutz // // 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_COPY_ON_DEVICE_HPP #define BOOST_COMPUTE_ALGORITHM_DETAIL_COPY_ON_DEVICE_HPP #include #include #include #include #include #include #include #include #include #include #include namespace boost { namespace compute { namespace detail { template inline event copy_on_device_cpu(InputIterator first, OutputIterator result, size_t count, command_queue &queue, const wait_list &events) { meta_kernel k("copy"); const device& device = queue.get_device(); 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" << "while(index < end){\n" << result[k.var("index")] << '=' << first[k.var("index")] << ";\n" << "index++;\n" << "}\n"; k.add_set_arg("count", static_cast(count)); size_t global_work_size = device.compute_units(); if(count <= 1024) global_work_size = 1; return k.exec_1d(queue, 0, global_work_size, events); } template inline event copy_on_device_gpu(InputIterator first, OutputIterator result, size_t count, command_queue &queue, const wait_list &events) { typedef typename std::iterator_traits::value_type input_type; const device& device = queue.get_device(); boost::shared_ptr parameters = detail::parameter_cache::get_global_cache(device); std::string cache_key = "__boost_copy_kernel_" + boost::lexical_cast(sizeof(input_type)); uint_ vpt = parameters->get(cache_key, "vpt", 4); uint_ tpb = parameters->get(cache_key, "tpb", 128); meta_kernel k("copy"); k << "uint index = get_local_id(0) + " << "(" << vpt * tpb << " * get_group_id(0));\n" << "for(uint i = 0; i < " << vpt << "; i++){\n" << " if(index < count){\n" << result[k.var("index")] << '=' << first[k.var("index")] << ";\n" << " index += " << tpb << ";\n" " }\n" "}\n"; k.add_set_arg("count", static_cast(count)); size_t global_work_size = calculate_work_size(count, vpt, tpb); return k.exec_1d(queue, 0, global_work_size, tpb, events); } template inline event dispatch_copy_on_device(InputIterator first, InputIterator last, OutputIterator result, command_queue &queue, const wait_list &events) { const size_t count = detail::iterator_range_size(first, last); if(count == 0){ // nothing to do return event(); } const device& device = queue.get_device(); // copy_on_device_cpu() does not work for CPU on Apple platform // due to bug in its compiler. // See https://github.com/boostorg/compute/pull/626 if((device.type() & device::cpu) && !is_apple_platform_device(device)) { return copy_on_device_cpu(first, result, count, queue, events); } return copy_on_device_gpu(first, result, count, queue, events); } template inline OutputIterator copy_on_device(InputIterator first, InputIterator last, OutputIterator result, command_queue &queue, const wait_list &events) { dispatch_copy_on_device(first, last, result, queue, events); return result + std::distance(first, last); } template inline discard_iterator copy_on_device(InputIterator first, InputIterator last, discard_iterator result, command_queue &queue, const wait_list &events) { (void) queue; (void) events; return result + std::distance(first, last); } template inline future copy_on_device_async(InputIterator first, InputIterator last, OutputIterator result, command_queue &queue, const wait_list &events) { event event_ = dispatch_copy_on_device(first, last, result, queue, events); return make_future(result + std::distance(first, last), event_); } #ifdef BOOST_COMPUTE_CL_VERSION_2_0 // copy_on_device() specialization for svm_ptr template inline svm_ptr copy_on_device(svm_ptr first, svm_ptr last, svm_ptr result, command_queue &queue, const wait_list &events) { size_t count = iterator_range_size(first, last); if(count == 0){ return result; } queue.enqueue_svm_memcpy( result.get(), first.get(), count * sizeof(T), events ); return result + count; } template inline future > copy_on_device_async(svm_ptr first, svm_ptr last, svm_ptr result, command_queue &queue, const wait_list &events) { size_t count = iterator_range_size(first, last); if(count == 0){ return future >(); } event event_ = queue.enqueue_svm_memcpy_async( result.get(), first.get(), count * sizeof(T), events ); return make_future(result + count, event_); } #endif // BOOST_COMPUTE_CL_VERSION_2_0 } // end detail namespace } // end compute namespace } // end boost namespace #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_COPY_ON_DEVICE_HPP