【发布时间】：2011-04-08 08:55:34
【问题描述】：
CUDA 支持递归吗？
【问题讨论】：
-
请注意，所有循环都涉及递归，而尾递归不/不应该涉及递归函数调用。
-
更多信息请见here
【发布时间】：2011-04-08 08:55:34
【问题描述】：
它适用于支持计算能力 2.0 和 CUDA 3.1 的 NVIDIA 硬件：
为 CUDA C 添加了新的语言功能 /C++ 包括：
支持功能 指针和递归使它更容易 将许多现有算法移植到 Fermi GPU
http://developer.nvidia.com/object/cuda_3_1_downloads.html
递归： 我在 NVIDIA 的网站上找不到代码示例，但在 forum 上有人发帖：
__device__ int fact(int f)
{
if (f == 0)
return 1;
else
return f * fact(f - 1);
}
【讨论】：
-
拥有“最新”硬件是不够的。并非所有最近的卡都是费米（又名计算能力 2.0）。目前，没有 Fermi 移动 GPU。
-
你说得对，我更新了我的 anwser。但是 Geforce GTX 480M 呢？里面有一个费米芯片。
-
有带费米的笔记本电脑； Geforce 480M 和 Quadro FX5000M 已经推出了一段时间。
-
确实如此。新的 GeForce GT 415M、420M、425M、435M、445M 和 GTX 460M、470M 即将到来！都具备计算能力 2.0。
-
我不认为这些是真正的递归调用，因为默认情况下 nvcc 会内联每个标记为 device 的函数。但是结果是一样的。
【讨论】：
即使它只支持特定芯片的递归，您有时也可以摆脱“模拟”递归：see how I used compile-time recursion for my CUDA raytracer。
【讨论】：
在 CUDA 4.1 版本中，CUDA 仅支持 __device__ 函数的递归，但不支持 __global__ 函数。
【讨论】：
仅在兼容设备上具有 2.0 计算能力之后
【讨论】：
当然可以，但它需要 Kepler 架构才能做到这一点。 查看他们关于经典快速排序的最新示例。
http://blogs.nvidia.com/2012/09/how-tesla-k20-speeds-up-quicksort-a-familiar-comp-sci-code/
据我所知，只有最新的 Kepler GK110 支持动态并行，它允许在内核中进行这种递归调用和产生新线程。在开普勒 GK110 之前，这是不可能的。请注意，并非所有的 Kepler 架构都支持这一点，只有 GK110 支持。
如果您需要递归，您可能需要 Tesla K20。 我不确定费米是否支持它，从来没有读过它。 :\ 但开普勒确实如此。 =)
【讨论】：
任何递归算法都可以用堆栈和循环来实现。这更痛苦，但如果你真的需要递归，这可以工作。
【讨论】：
CUDA 3.1 支持递归
【讨论】：
如果您的算法涉及大量递归，那么无论是否支持，它都不是为 GPU 设计的，要么重新设计您的算法，要么获得更好的 CPU，无论哪种方式都会更好（我敢打赌，在很多情况下，数量级会更好）然后在 GPU 上进行递归。
【讨论】：
是的，实际版本支持。但是尽管可以执行递归函数，但您必须记住，无法预测执行堆栈中的内存分配（必须执行递归函数才能知道递归的真实深度），所以您的堆栈可能导致无法满足您的目的，并且可能需要手动增加默认堆栈大小
【讨论】：
是的，它确实支持递归。但是，在 GPU 上进行递归并不是一个好主意。因为每个线程都会这样做。
【讨论】：
-
引用（文档等）将使这个答案更完整。作为参考，似乎添加了 in CUDA 3.1：“添加到 CUDA C/C++ 的新语言特性包括：对函数指针和递归的支持使得将许多现有算法移植到 Fermi GPU 变得更加容易”
-
没错。此功能已添加到工具包 3.1 中的 CUDA C。最新版本的 CUDA 编程指南暗示支持递归设备功能。但是
__global__函数不支持递归。请参阅 F.3.9.6。编程指南中的函数递归。
刚刚在我的电脑上使用具有 1.1 计算能力的 NVIDIA GPU 进行了尝试。它说尚不支持递归。所以它与运行时无关，而是与硬件本身有关
【讨论】：