inline void initial_reduce(InputIterator first,
InputIterator last,
buffer result,
const Function &function,
kernel &reduce_kernel,
const uint_ vpt,
const uint_ tpb,
command_queue &queue)
{
(void) function;
(void) reduce_kernel;
typedef typename std::iterator_traits<InputIterator>::value_type Arg;
typedef typename boost::tr1_result_of<Function(Arg, Arg)>::type T;
size_t count = std::distance(first, last);
detail::meta_kernel k("initial_reduce");
k.add_set_arg<const uint_>("count", uint_(count));
size_t output_arg = k.add_arg<T *>(memory_object::global_memory, "output");
k <<
k.decl<const uint_>("offset") << " = get_group_id(0) * VPT * TPB;\n" <<
k.decl<const uint_>("lid") << " = get_local_id(0);\n" <<
"__local " << type_name<T>() << " scratch[TPB];\n" <<
// private reduction
k.decl<T>("sum") << " = 0;\n" <<
"for(uint i = 0; i < VPT; i++){\n" <<
" if(offset + lid + i*TPB < count){\n" <<
" sum = sum + " << first[k.var<uint_>("offset+lid+i*TPB")] << ";\n" <<
" }\n" <<
"}\n" <<
"scratch[lid] = sum;\n" <<
// local reduction
ReduceBody<T,false>::body() <<
// write sum to output
"if(lid == 0){\n" <<
" output[get_group_id(0)] = scratch[0];\n" <<
"}\n";
const context &context = queue.get_context();
std::stringstream options;
options << "-DVPT=" << vpt << " -DTPB=" << tpb;
kernel generic_reduce_kernel = k.compile(context, options.str());
generic_reduce_kernel.set_arg(output_arg, result);
size_t work_size = calculate_work_size(count, vpt, tpb);
queue.enqueue_1d_range_kernel(generic_reduce_kernel, 0, work_size, tpb);
}
event exec(command_queue &queue)
{
if(m_count == 0){
// nothing to do
return event();
}
size_t global_work_size = calculate_work_size(m_count, m_vpt, m_tpb);
set_arg(m_count_arg, uint_(m_count));
return exec_1d(queue, 0, global_work_size, m_tpb);
}
inline void initial_reduce(const buffer_iterator<T> &first,
const buffer_iterator<T> &last,
const buffer &result,
const plus<T> &function,
kernel &reduce_kernel,
const uint_ vpt,
const uint_ tpb,
command_queue &queue)
{
(void) function;
size_t count = std::distance(first, last);
reduce_kernel.set_arg(0, first.get_buffer());
reduce_kernel.set_arg(1, uint_(first.get_index()));
reduce_kernel.set_arg(2, uint_(count));
reduce_kernel.set_arg(3, result);
reduce_kernel.set_arg(4, uint_(0));
size_t work_size = calculate_work_size(count, vpt, tpb);
queue.enqueue_1d_range_kernel(reduce_kernel, 0, work_size, tpb);
}
