Is there any way to stop OpenCL kernel from execution?

Question

Is there any way to stop OpenCL kernel from execution? For example, I launch the kernel, do some computations, and then stop it if some conditions were met, otherwise, I wait until it finishes:

clEnqueueNDRange(queue, ...); // start kernel function

// do other stuff...
// ...

if (some condition met) {
    stopKernel();
} else { 
    clFinish(queue);
}

Thank you for help

Answer 1

No. Once you have enqueued your kernel, it will run to completion.

One way to accomplish something like the above is to do this:

while ( data_left_to_process ) {

   clEnqueueNDRangeKernel( ..., NDRange for a portion of the data, ... )

   // other work

   if (condition) {
      break;
   }

   // adjust NDRange for next execution to processes the next part of your data

}

clFinish(queue);

This allows you avoid processing ALL the data, with the obvious tradeoff that you're now submitting work in smaller chunks, which will probably have a performance impact.

Answer 2

Possibly.

1. Create two command queues in a context.
2. Create two kernels, one to do the work and another to halt execution. Each kernel has access to a shared global buffer.
3. Load the first kernel into queue1.
4. Load the second kernel into queue2 when you want to halt execution.

Alternatively you could use an out-of-order queue and load the second kernel into the the same command queue to halt execution. You have to be a little more careful (using clFinish/clFlush as necessary), however it is a more natural way of doing this.

Some pseudo code (for multiple queues):

clEnqueueNDRange(queue1, kernel1, ...); //start work kernel
// do other stuff
// ...
if (condition)
    clEnqueueNDRange(queue2, kernel2, ...); //stop work kernel
clFinish(queue1);
clFinish(queue2); // entirely unnecessary if using in-order queues

Use a buffer of ints or floats as your stopping variable and access them via the global_id within your kernels to reduce the cost of reading from global within a loop. The downside is that your state will be indeterminate: without further variables to count executions etc, you won't know how many work items and which ones have been executed.

And the kernels:

void kernel1( ... ,global int * g_stop)
{
    int index_from_ids = ...;
    while (g_stop[index_from_ids] == 0) // or 'if' for single pass execution
    {
        // do work here
    }
}

void kernel2( ... ,global int * g_stop)
{
    int index_from_ids = ...;
    g_stop[index_from_ids] = 1;
}