How to cancel a running cuda kernel?

Is there anyway to cancel a cuda kernel after it has started execution? The goal here is to free gpu resources as soon as possible for other uses.

Thanks in advance for any suggestions.

From inside a kernel you can use [font=“Courier New”]asm(“trap;”)[/font] for that.

Thnx tera for your reply. Is asm(“trap;”) faster than a simple return? Also, how can i inform cuda to abort execution? I was thinking of a flag in global mapped memory that could be written by the host and then be read by cuda thread blocks, but is there any faster way?

[font=“Courier New”]asm(“trap;”)[/font] aborts all threads, while [font=“Courier New”]return[/font] (or [font=“Courier New”]asm(“exit;”)[/font]) just ends the current thread. Also [font=“Courier New”]asm(“trap;”)[/font] will signal an error to the host (general launch failure IIRC, while the others return success.

The flag in global memory is the way to go. I don’t know of any faster way. Don’t put the flag in mapped memory though, or each thread/block would have to read it through PCIe though. You can use cudaMemcpy() to set the flag from the host.

Hi guys,

tera answer me in another topic, same as this topic. I have the same problem, i try to initialize a flag variable on the host (or off the host in a variable device), and put the variable in the kernel via cudaMemcpy(), but i don’t know if the flag is see for all the threads and when i flag change his value, the others threads notice that flag is changed and stop the execution.

My topic: http://forums.nvidia.com/index.php?showtopic=222058&st=0&gopid=1363744&#entry1363744