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

//---------------------------------------------------------------------------// // 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_INPLACE_REDUCE_HPP #define BOOST_COMPUTE_ALGORITHM_DETAIL_INPLACE_REDUCE_HPP #include #include #include #include #include #include namespace boost { namespace compute { namespace detail { template inline void inplace_reduce(Iterator first, Iterator last, BinaryFunction function, command_queue &queue) { typedef typename std::iterator_traits::value_type value_type; size_t input_size = iterator_range_size(first, last); if(input_size < 2){ return; } const context &context = queue.get_context(); size_t block_size = 64; size_t values_per_thread = 8; size_t block_count = input_size / (block_size * values_per_thread); if(block_count * block_size * values_per_thread != input_size) block_count++; vector output(block_count, context); meta_kernel k("inplace_reduce"); size_t input_arg = k.add_arg(memory_object::global_memory, "input"); size_t input_size_arg = k.add_arg("input_size"); size_t output_arg = k.add_arg(memory_object::global_memory, "output"); size_t scratch_arg = k.add_arg(memory_object::local_memory, "scratch"); k << "const uint gid = get_global_id(0);\n" << "const uint lid = get_local_id(0);\n" << "const uint values_per_thread =\n" << uint_(values_per_thread) << ";\n" << // thread reduce "const uint index = gid * values_per_thread;\n" << "if(index < input_size){\n" << k.decl("sum") << " = input[index];\n" << "for(uint i = 1;\n" << "i < values_per_thread && (index + i) < input_size;\n" << "i++){\n" << " sum = " << function(k.var("sum"), k.var("input[index+i]")) << ";\n" << "}\n" << "scratch[lid] = sum;\n" << "}\n" << // local reduce "for(uint i = 1; i < get_local_size(0); i <<= 1){\n" << " barrier(CLK_LOCAL_MEM_FENCE);\n" << " uint mask = (i << 1) - 1;\n" << " uint next_index = (gid + i) * values_per_thread;\n" " if((lid & mask) == 0 && next_index < input_size){\n" << " scratch[lid] = " << function(k.var("scratch[lid]"), k.var("scratch[lid+i]")) << ";\n" << " }\n" << "}\n" << // write output for block "if(lid == 0){\n" << " output[get_group_id(0)] = scratch[0];\n" << "}\n" ; const buffer *input_buffer = &first.get_buffer(); const buffer *output_buffer = &output.get_buffer(); kernel kernel = k.compile(context); while(input_size > 1){ kernel.set_arg(input_arg, *input_buffer); kernel.set_arg(input_size_arg, static_cast(input_size)); kernel.set_arg(output_arg, *output_buffer); kernel.set_arg(scratch_arg, local_buffer(block_size)); queue.enqueue_1d_range_kernel(kernel, 0, block_count * block_size, block_size); input_size = static_cast( std::ceil(float(input_size) / (block_size * values_per_thread) ) ); block_count = input_size / (block_size * values_per_thread); if(block_count * block_size * values_per_thread != input_size) block_count++; std::swap(input_buffer, output_buffer); } if(input_buffer != &first.get_buffer()){ ::boost::compute::copy(output.begin(), output.begin() + 1, first, queue); } } } // end detail namespace } // end compute namespace } // end boost namespace #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_INPLACE_REDUCE_HPP