Does CUDA support recursion?

Question

Does CUDA support recursion?

Answer 1

This applies to NVIDIA hardware that supports computing capabilities 2.0 and CUDA 3.1:

New language features added in CUDA C / C ++ include :

Function support pointers and recursion make it easy to port many existing algorithms to Fermi GPUs

http://developer.nvidia.com/object/cuda_3_1_downloads.html

Function pointers: http://developer.download.nvidia.com/compute/cuda/sdk/website/CUDA_Advanced_Topics.html#FunctionPointers

recursion: I can’t find the sample code on the NVIDIA website, but someone posted this message on the forum :

__device__ int fact(int f) { if (f == 0) return 1; else return f * fact(f - 1); } 

Answer 2

Yes, see NVIDIA CUDA Programming Guide :

device functions only support recursion in device code compiled for device 2.0 computing capabilities.

You need a Fermi map to use them.

Answer 3

Despite the fact that it only supports recursion for certain chips, you can sometimes get away with "emulated" recursion: see how I used recursion at compile time for my CUDA tracer .

Answer 4

In the release of CUDA 4.1, CUDA only supports recursion for the __device__ function, but not for the __global__ function.

Answer 5

Only after the possibility of computing 2.0 on compatible devices

Answer 6

Of course, this is true, but this requires the Kepler architecture. Check out their latest classic quicksort example.

http://blogs.nvidia.com/2012/09/how-tesla-k20-speeds-up-quicksort-a-familiar-comp-sci-code/

As far as I know, only the latest Kepler GK110 supports dynamic parallelism, which allows this kind of recursive call and spawning of new threads in the kernel. Prior to the Kepler GK110, this was not possible. And note that not all Kepler architecture supports this, only the GK110 does.

If you need recursion, you probably need a Tesla K20. I'm not sure if Fermi supports, I never read about it .: \ But Kepler is sure. =)

Answer 7

Any recursive algorithm can be implemented using a stack and a loop. This is more of a pain, but if you really need recursion, it might work.

Answer 8

CUDA 3.1 supports recursion

Answer 9

If your algorithm causes many recursions, then support or not, it is not intended for GPUs, or redesigns of your algorti, or to get a better processor, since it will be better (I put, in many cases, the magic is better), then do the second one GPU verification.

Answer 10

Yes, it is supported in the real version. But, despite the fact that you can perform recursive functions, you should keep in mind that memory allocation from the execution stack cannot be predicted (a recursive function must be executed to find out the true depth of the recursion), so your stack may not be sufficient for your purposes, and he may need to manually increase the default stack size

Answer 11

Yes, it supports recursion. However, recursing on a GPU is not a good idea. Because every thread is going to do it.

Answer 12

I tried right now on my computer with an NVIDIA GPU with 1.1 Compute capabilities. He says recursion is not yet supported. Thus, it has nothing to do with runtime, but the device itself