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

Question

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

例如，以下代码片段使用相同数量的寄存器吗？他们的效率相同吗？

片段 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);
}

Answer 1

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

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

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

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

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

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

1. 如果该函数位于与调用它的函数或其他函数__device__不同的编译单元中。在这种情况下，唯一可行的方法是通过CUDA 单独编译和链接，也称为设备链接。在这种情况下，编译器将不会（无法）内联该函数。__global____device__

2. 如果指定了__noinline__ 编译器指令。请注意，这只是对编译器的提示；它可能会被忽略。