PDA

View Full Version : reg kernel function



paviraj_1989
03-31-2011, 08:37 AM
Suppose in Kernel function there will be n work items performing some task, if one of the work item finishes its task early then it has to terminate and all other work item must also terminate. Can you help me how to do this

david.garcia
03-31-2011, 04:13 PM
There's no good way to do that. You could use something like a __local integer variable and force work-items to poll the value of that variable. When a work-item finishes it will set the variable to 1, and when other work-items find that the variable is set to 1 they will return.

The kernel would look something like this:


__kernel void foo()
{
__local int done;

// The first work-item sets 'done=0'
if(!get_global_id(0))
{
done = 0;
}

// Ensure all work items see 'done=0' before they start to do any work.
barrier(CLK_LOCAL_MEM_FENCE);

for(...)
{
// Do work here...
if(this_work_item_found_the_solution)
{
done = 1;
}

// All work items check whether the solution has been found
barrier(CLK_LOCAL_MEM_FENCE);
if(done)
{
return;
}
}
}



There may be better ways to do it. You could use atomics instead of barriers, although I don't know which one will be slower (it will depend on the implementation).