Cuda 中的线程内如何管理堆栈帧？

Question

假设我们有一个调用一些函数的内核，例如：

__device__ int fib(int n) {
    if (n == 0 || n == 1) {
        return n;
    } else {
        int x = fib(n-1);
        int y = fib(n-2);
        return x + y;
    }
    return -1;
}

__global__ void fib_kernel(int* n, int *ret) {
    *ret = fib(*n);
}

内核fib_kernel将调用 function fib()，该函数在内部将调用两个fib()函数。假设GPU有80个SM，我们正好启动80个线程来进行计算，并传入n10个。我知道会有大量的重复计算，这违反了数据并行性的思想，但我想更好地理解线程的堆栈管理。

根据Cuda PTX 的文档，它声明如下：

GPU 维护每个线程的执行状态，包括程序计数器和调用堆栈

1. 堆栈位于本地内存中。当线程执行内核时，它们的行为是否与CPU中的调用约定一样？换句话说，是不是对于每个线程，对应的栈都会动态增长和收缩呢？

2. 每个线程的堆栈都是私有的，其他线程无法访问。有没有一种方法可以手动检测编译器/驱动程序，以便堆栈分配在全局内存中，而不是本地内存中？

3. 有没有一种方法可以让线程获取当前的程序计数器、帧指针值？我认为它们存储在一些特定的寄存器中，但 PTX 文档没有提供访问这些寄存器的方法。我可以知道我必须修改什么（例如驱动程序或编译器）才能获取这些寄存器吗？

4. 如果我们将输入增加到fib(n)10000，很可能会导致堆栈溢出，有办法处理吗？问题2的答案或许可以解决这个问题。任何其他想法将不胜感激。

Answer 1

如果您从几个示例中研究生成的 SASS 代码，您将会更好地了解这些事情是如何工作的。

当线程执行内核时，它们的行为是否与CPU中的调用约定一样？换句话说，是不是对于每个线程，对应的栈都会动态增长和收缩呢？

CUDA 编译器将尽可能积极地内联函数。当它不能时，它会在本地内存中构建一个类似堆栈的结构。然而，我知道的 GPU 指令不包括显式堆栈管理（例如，推入和弹出），因此“堆栈”是“由编译器构建”，使用保存（本地）地址的寄存器LD/ST 指令将数据移入/移出“堆栈”空间。从这个意义上说，实际堆栈的大小确实/可以动态改变，但是最大允许的堆栈空间是有限的。每个线程都有自己的堆栈，使用此处给出的“堆栈”的定义。

有没有一种方法可以手动检测编译器/驱动程序，以便堆栈分配在全局内存中，而不是本地内存中？

实际上，不。生成指令的 NVIDIA 编译器有一个前端和一个闭源的后端。如果你想修改 GPU 的开源编译器，这是可能的，但目前据我所知，没有广泛认可的工具链不使用闭源后端（ptxas或其等效驱动程序） ）。GPU 驱动程序也是大型闭源的。也没有任何会影响堆栈位置的公开控件。

我可以知道我必须修改什么（例如驱动程序或编译器）才能获取这些寄存器吗？

指令指针/程序计数器没有公开的寄存器。因此不可能说明需要进行哪些修改。

如果我们将fib(n)的输入增加到10000，很可能会导致堆栈溢出，有办法处理吗？

正如我所提到的，每个线程的最大堆栈空间是有限的，因此您的观察是正确的，最终堆栈可能会增长到超过可用空间（这可能是 CUDA 设备代码中递归的危险）。提供的解决此问题的机制是增加每线程本地内存大小（因为堆栈存在于逻辑本地空间中）。