1

在 CUDA 5 编程指南中，如下所述：

发布可能会持续到 24 代的深度，但这个深度通常会受到 GPU 上可用资源的限制

我的问题如下：

  • GPU 上的 CUDA 运行时是否保证始终可以达到 24 的深度，并且在某些情况下甚至可能超过 24（案例 A）？还是他们的意思是 24 是绝对最大限制，并且在运行时可能确实无法达到这个数字（案例 B）？

  • 如果是情况 B，当内核在 GPU 上启动并且没有足够的资源时会发生什么？发射失败了吗？（如果是这种情况很奇怪！）

我计划编写一个 CUDA 程序，我想从 Kepler 架构中受益。我的算法通常绝对需要 15-19 级别的函数递归（递归级别绑定到我的数据结构）。

参考：TechBrief_Dynamic_Parallelism_in_CUDA_v2.pdf

4

1 回答 1

1

CUDA 不保证会实现递归深度 1 - 与传统操作系统不保证启动新进程/线程会成功类似。例如，如果您有以下程序：

int main() {
    pid_t pid;
    while (pid = fork ());
    while (true) {
        dummy<<<1, 1024>>> ();
    }
}

__global void dummy() {}

在某些时候某些事情会失败——要么你将耗尽 CPU 或 GPU 内存。在 GPU 上以类似的方式，它可能会失败（返回错误 - CUDA 或 fork 将返回 -1）。

另一种看待它的方式是，在最坏的情况下，每次启动都可以有(2^31-1)^2*(2^10-1) ≃ 2^72每个带有线程的块。2^10即在单次启动中你可以有2^82线程。现在每个递归都是指数级的，因此即使在最坏的情况下启动后终止线程，它也需要保证调度2^1968线程。如果每个线程的状态是 1/32 位，无论是否完成扭曲，它都需要2^1945GiB 的内存（这“略”多于观察到的宇宙的信息容量 - 即2^1595多倍）。

因此，它肯定是案例 B，并且不可能是案例 A（扭曲状态必须至少包括指令指针）。根据分支因子，如果您同步，您的算法可能可以实现 15-19 的递归深度。

编辑：如果您的意思是普通递归而不是递归启动，那么实际上它受到堆栈的限制。根据确切的代码，它在 Fermi+ 上实际上可能是无限的（特斯拉一代不支持递归 IIRC）。同样，没有保证的最小深度——在堆栈/本地内存上分配大数组，你会用完空间（优化器很好地摆脱了它）。

于 2013-01-13T11:07:26.247 回答