kernel abort/return value

Accelerated Computing CUDA CUDA Programming and Performance
1

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?

Thanks

2

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.

eyal

3

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…

4

Nope, no other way.

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

eyal