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

Question

我已经阅读过各种地方，__device__函数几乎总是由CUDA编译器内联。那么说，当我将代码从内核移动到由内核调用的__device__函数中时，所使用的寄存器数量（通常）不会增加？

作为一个例子，下面的代码段使用相同数量的寄存器吗？他们有同样的效率吗？

SNIPPET 1

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

SNIPPET 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__函数可以影响所用寄存器的数量（如w性能和效率）。

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

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

1. 它必须从堆栈框架移动由所述函数中使用的变量来/
2. 基于“周围”的代码，不能进行其他优化，因为它不知道周围的代码是什么。编译器必须以独立的方式处理__device__函数。

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

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

1. 如果__device__功能是在一个单独的编译单元从调用它的__global__或其他__device__功能。在这种情况下，唯一可行的方法是通过CUDA separate compilation and linking，也称为设备链接。在这种情况下，编译器不会（不能）内联该函数。

2. 如果指定__noinline__compiler directive。请注意，这只是编译器的一个提示;它可能被忽略。