kernel abort/return value

I know this may look tricky, but it would be really useful in some cases.
Is there any way I can abort a kerner execution from inside the kernel and return some value?

Sometimes, I have a lot of data to process, but if I find an error, I don’t want to continue kernel execution and I want to return as fast as possible to the host code.

Let’s suppose the following simple example. I want to check that every element in a vector with 10 million values is >= 5.
As soon as I find something below 5, I want to abort and return false. I don’t want execute all blocks/grid, etc… Just abort and return some value that indicate something like UserAbortCode=1.

Any ideas?


You can define some int value in gmem, initialize it to zero. in case you find some value that is less than 5 in any

thread/block you set it to 1 (if i recall correctly even if many threads or blocks update the same value to 1, you should

still see one).

Everytime a block starts have threadIdx.x == 0 check this value, if its zero do the calculation as above, otherwise return.

There is no special function to quit the kernel.


Yes, we were using something like that… but then I access global memory a lot of times, and I still have to run all the blocks.

Are you 100% sure I can’t somehow abort a kernel? I would love to be able to do something like that…

Nope, no other way.

Make sure you only read this flag once per block thread and not per thread.