Does CUDA support recursion?

Solution 1:

It does on NVIDIA hardware supporting compute capability 2.0 and CUDA 3.1:

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

Support for function pointers and recursion make it easier 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 a code sample on NVIDIA's website, but on the forum someone post this:

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

Solution 2:

Yes, see the NVIDIA CUDA Programming Guide:

device functions only support recursion in device code compiled for devices of compute capability 2.0.

You need a Fermi card to use them.

Solution 3:

Even though it only supports recursion for specific chips, you can sometimes get away with "emulated" recursion: see how I used compile-time recursion for my CUDA raytracer.

Solution 4:

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

Solution 5:

Only after 2.0 compute capability on compatible devices