调用 __device__ 函数会影响 CUDA 中使用的寄存器数量吗？答案

【问题标题】：Does calling __device__ functions impact the number of registers used in CUDA?调用 __device__ 函数会影响 CUDA 中使用的寄存器数量吗？
【发布时间】：2016-03-06 21:36:06
【问题描述】：

我在不同的地方读到 __device__ 函数几乎总是被 CUDA 编译器内联。那么，当我将代码从内核移动到内核调用的__device__ 函数时，（通常）使用的寄存器数量没有增加是否正确？

举个例子，下面的sn-ps是否使用相同数量的寄存器？它们是否同样有效？

片段 1

__global__ void manuallyInlined(float *A,float *B,float *C,float *D,float *E) {
    // code that manipulates A,B,C,D and E 
}

片段 2

__device__ void fn(float *A,float *B,float *C,float *D,float *E) {
    // code that manipulates A,B,C,D and E 
}


__global__ void manuallyInlined(float *A,float *B,float *C,float *D,float *E) {
    fn(A,B,C,D,E);
}

【问题讨论】：

标签： performance cuda inline

【解决方案1】：

最终答案只能通过使用工具来确定（使用-Xptxas -v 编译，或使用其中一个分析器），但一般的答案是调用__device__ 函数可以影响使用的寄存器数量（以及性能和效率）。

根据您的文件组织以及编译代码的方式，__device__ 函数可能是 inlined。如果它是内联的，这通常会为优化编译器（主要是ptxas）提供最佳机会来适应它认为合适的寄存器使用。（请注意，至少在理论上，这种“适应”可能会导致 either 使用更多或更少的寄存器。但是，内联情况通常会导致编译器使用更少的寄存器和可能更高的性能。但是编译器主要优化了更高的性能，而不是更少的寄存器使用。）

另一方面，如果它不是内联的，那么它必须作为一个普通的函数调用来处理。与许多其他计算机体系结构一样，函数调用涉及设置堆栈帧以传递变量，然后将控制权转移给函数。在这种情况下，编译器受到更多限制，因为：

它必须将函数使用的变量移入/移出堆栈帧
它不能基于“周围”代码执行其他优化，因为它不知道周围的代码是什么。 __device__ 函数必须由编译器以独立方式处理。

因此，如果函数可以内联，那么您的两种方法之间应该没有太大区别。如果函数不能内联，那么上述两种方法的寄存器使用通常会有明显的差异。

可能会影响编译器是否会尝试内联 __device__ 函数的一些明显因素是：

如果 __device__ 函数与调用它的 __global__ 或其他 __device__ 函数位于单独的编译单元中。在这种情况下，唯一可行的方法是通过CUDA separate compilation and linking，也称为设备链接。在这种情况下，编译器不会（不能）内联函数。
如果指定了__noinline__ compiler directive。请注意，这只是对编译器的提示；可以忽略。

【讨论】：