CODEWITHSUNDEEP
recursioncuda
JuanPablo
JuanPablo

Reputation: 24764

Does CUDA support recursion?

Does CUDA support recursion?

Upvotes: 63

Views: 32904

Answers (12)

palebluedot
palebluedot

Reputation: 15

Yes, it does support recursion. However, it is not a good idea to do recursion on GPU. Because each thread is going to do it.

Upvotes: 0

Mr.PotatusVII
Mr.PotatusVII

Reputation: 48

Yeah, it is supported on the actual version. But despite the fact it is possible to execute recursive functions, you must have in mind that the memory allocation from the execution stack cannot be predicted (the recursive function must be executed in order to know the true depth of the recursion), so your stack could result being not enough for your purposes and it could need a manual increment of the default stack size

Upvotes: 0

Arturo Garcia
Arturo Garcia

Reputation: 51

Only after 2.0 compute capability on compatible devices

Upvotes: 5

user0002128
user0002128

Reputation: 2921

If your algorithm invovles alot of recursions, then support or not, it is not designed for GPUs, either redesign your algorthims or get a better CPU, either way it will be better (I bet in many cases, maginitudes better) then do recurisons on GPUs.

Upvotes: 1

Hong Zhou
Hong Zhou

Reputation: 649

Sure it does, but it requires the Kepler architecture to do so. Check out their latest example on the classic quick sort.

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

As far as i know, only latest Kepler GK110 supports dynamic parallelism, which allow this kind of recursive call and spawning of new threads within the kernel. Before Kepler GK110, it was not possible. And note that not all Kepler architecture supports this, only GK110 does.

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

Upvotes: 3

username_4567
username_4567

Reputation: 4903

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

Upvotes: 7

sp497
sp497

Reputation: 2491

Tried just now on my pc with a NVIDIA GPU with 1.1 Compute capability. It says recursion not yet supported. So its not got anything to do with the runtime but the hardware itself

Upvotes: -2

ttsiodras
ttsiodras

Reputation: 11258

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.

Upvotes: 9

elmattic
elmattic

Reputation: 12174

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);
}

Upvotes: 54

Jan C
Jan C

Reputation: 21

CUDA 3.1 supports recursion

Upvotes: 2

Dr. Snoopy
Dr. Snoopy

Reputation: 56357

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.

Upvotes: 13

dicroce
dicroce

Reputation: 46770

Any recursive algorithm can be implemented with a stack and a loop. It's way more of a pain, but if you really need recursion, this can work.

Upvotes: 3

Related Questions