为什么在 CUDA 中使用递归时会收到警告？

Question

我必须使用CUDA解决 GPU 上的代码问题，但我总是收到堆栈大小的警告，因为“函数名称”无法静态确定。

这是我正在开发的一个学生项目，该项目是使用CUDA 9.0库用C语言编写的，并且在NVIDIA Quadro K5000 gpu上运行。 每个线程必须执行一个函数，并且在这个函数中，同一个函数有两个递归调用，我之所以想使用这两个递归调用，是因为它使代码对我来说干净简单，但是如果有只有一次递归调用，不再存在堆栈大小问题。

这是我每次编译代码时遇到的错误：

CUDA支持递归函数调用，但我不明白为什么当有两个递归调用时会出现问题。

__device__ void bitonicMergeGPU(float *arr, int l, int indexT, int order) 
{
    int k,p;

    if(l > 1)
    {
        p = l/2;

        for(k=indexT;k<indexT+p;k++)
        {
            //Compare the values.
            compareAndExchange(arr,k,k+p,order);
        }

        //THIS IS WHERE I GET THE ERROR

        bitonicMergeGPU(arr,p,indexT,order);
        bitonicMergeGPU(arr,p,indexT+p,order);
    }
}

我只是想知道是否可以解决递归调用的问题。

Answer 1

CUDA 支持递归。当您在 CUDA 中使用递归时，会出现此警告，并且没有 NVIDIA 记录的方法可以使警告消失（除非不使用递归）。

如果递归地使用函数，在大多数语言中，随着递归深度的增加，它将使用更多的堆栈空间。在 CUDA 中也是如此。您需要考虑到这一点，并为您预期的最大递归深度提供足够的堆栈空间。限制递归深度是常见的做法，以防止堆栈问题。

编译器无法在编译时发现最大运行时递归深度，并且会出现警告来提醒您这一点。

无论您增加多少堆栈大小，警告都不会消失。该警告是为了让您知道您有责任确保您的递归设计以及分配的堆栈空间能够正常工作。编译器不会以任何方式验证堆栈大小的增加量是否足够。