//-----------------------------------------------------------------------------
void kernel::execute( const generic_command_queue* queue_ptr,
const ndrange& offset, const ndrange& global,
const ndrange& local, events_t& wait_events,
generic_event** event_ptr )
{
const command_queue* queue = reinterpret_cast<const command_queue*>( queue_ptr );
cl_int error_code;
cl_event evnt;
if( wait_events.empty() )
{
error_code =
opencl_.clEnqueueNDRangeKernel( queue->get_id(), get_id(),
static_cast<cl_uint>( global.get_dimensions() ),
offset.get_pointer(), global.get_pointer(),
local.get_pointer(), 0, NULL, &evnt );
}
else
{
boost::scoped_array<cl_event> events( new cl_event[wait_events.size()] );
for( uint32_t i = 0; i < wait_events.size(); i++ )
{
events[ i ] = (reinterpret_cast<const event*>( wait_events[ i ] ))->get_id();
}
error_code =
opencl_.clEnqueueNDRangeKernel( queue->get_id(), get_id(),
static_cast<cl_uint>( global.get_dimensions() ),
offset.get_pointer(), global.get_pointer(), local.get_pointer(),
static_cast<cl_uint>( wait_events.size() ), events.get(), &evnt );
}
if( error_code != CL_SUCCESS )
{
throw dcl::library_exception( error_code );
}
if( (event_ptr != NULL) && (evnt != NULL) )
{
*event_ptr = new event( opencl_, queue_ptr, evnt );
}
else if( evnt != NULL )
{
opencl_.clReleaseEvent( evnt );
}
}
Event count(const events_t & events) {
if (events.empty()) {
return {};
}
return events.front().copy().set_metric(events.size())
.set_time(max_time(events));
}
Event mean(const events_t & events) {
if (events.empty()) {
return {};
}
Event e(events.front());
e.set_metric_d(reduce(sum_fn, events) / events.size());
e.set_time(max_time(events));
return e;
}