CUDA __global__ 函数递归答案

【问题标题】：CUDA __global__ function recursionCUDA __global__ 函数递归
【发布时间】：2013-09-19 20:27:40
【问题描述】：

我在 CUDA 中有一个 __global__ 函数。它可以调用自己吗？

这是我的例子：

__global__ void 
force_create_empty_nodes (struct NODE *Nodes, int topnode, int bits, int no, int x, int y, 
                          int z, struct topnode_data *TopNodes)
{
    /// * Some code *///
    force_create_empty_nodes <<<1, 8>>>(Nodes, topnode+1, bits+1, no+1, 
                                             x+1, y+1, z+1, TopNodes);
}

我收到的错误是：

error: kernel launch from __device__ or __global__ functions requires separate compilation mode

这是我的 make 命令：

nvcc -c -arch compute_35 cudaForceNodes.cu -o obj/cudaForceNodes.o

【问题讨论】：

看看这个问题的答案：stackoverflow.com/questions/3644809/does-cuda-support-recursion
我看过这里，但我的 CUDA 版本是 5.5，我的 GPU 是 GeForce GT 550M。我尝试了 device 功能并且可以工作，但是尝试 global 功能时我无法构建项目。我需要在每次递归中创建新线程。是否可以通过其他方式实现。
如果您拥有 GT 550M，则不支持动态并行（即从内核启动内核），并且为 sm_25 架构编译的代码将无法在您的 GPU 上运行。

标签： cuda

【解决方案1】：

从另一个内核调用一个内核称为动态并行。它的文档是here。

需要：

计算能力 3.5 设备。您可以通过运行 cuda deviceQuery 示例找到您设备的计算能力。
编译命令中的各种开关，包括为 cc3.5 架构指定编译和单独（设备）编译所需的开关，以及与设备运行时链接的开关。

由于您的 GT550M 不是 cc 3.5 设备，您将无法使用此功能。没有其他方法可以从内核中调用内核。

【讨论】：