调用 __device__ 函数会影响 CUDA 中使用的寄存器数量吗?
Does calling __device__ functions impact the number of registers used in CUDA?
我在各个地方读到 __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);
}
最终答案只能通过使用工具来确定(使用 -Xptxas -v
编译或使用其中一个分析器),但一般答案是调用 __device__
函数 可以影响使用的寄存器数量(以及性能和效率)。
根据您的文件组织以及您编译代码的方式,__device__
函数可能是 inlined. If it is inlined, this generally gives the optimizing compiler (ptxas,主要是)在它认为合适的情况下调整寄存器使用的最佳机会。 (请注意,至少在理论上,此 "adaptation" 可能会导致 或者 使用更多或更少的寄存器。但是,内联情况通常会导致编译器使用更少的寄存器和可能更高的性能。但是编译器主要优化以获得更高的性能,而不是更少的寄存器使用。)
另一方面,如果它不是内联的,那么它必须作为一个普通的函数调用来处理。与许多其他计算机体系结构一样,函数调用涉及设置堆栈帧以传递变量,然后将控制权转移给函数。在这种情况下,编译器受到更多限制,因为:
- 它必须将函数to/from使用的变量移动到栈帧
- 它无法在"surrounding"代码的基础上进行其他优化,因为它不知道周围的代码是什么。
__device__
函数必须由编译器以独立方式处理。
因此,如果函数可以内联,那么您的两种方法之间应该没有太大区别。如果函数不能内联,那么上述两种方式在寄存器使用上通常会有明显的区别。
可能影响编译器是否尝试内联 __device__
函数的一些明显因素是:
如果 __device__
函数与调用它的 __global__
或其他 __device__
函数位于单独的编译单元中。在这种情况下,唯一可行的方法是通过 CUDA separate compilation and linking,也称为设备链接。在这种情况下,编译器将不会(不能)内联函数。
如果指定了__noinline__
compiler directive。请注意,这只是对编译器的提示;它可能会被忽略。
我在各个地方读到 __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);
}
最终答案只能通过使用工具来确定(使用 -Xptxas -v
编译或使用其中一个分析器),但一般答案是调用 __device__
函数 可以影响使用的寄存器数量(以及性能和效率)。
根据您的文件组织以及您编译代码的方式,__device__
函数可能是 inlined. If it is inlined, this generally gives the optimizing compiler (ptxas,主要是)在它认为合适的情况下调整寄存器使用的最佳机会。 (请注意,至少在理论上,此 "adaptation" 可能会导致 或者 使用更多或更少的寄存器。但是,内联情况通常会导致编译器使用更少的寄存器和可能更高的性能。但是编译器主要优化以获得更高的性能,而不是更少的寄存器使用。)
另一方面,如果它不是内联的,那么它必须作为一个普通的函数调用来处理。与许多其他计算机体系结构一样,函数调用涉及设置堆栈帧以传递变量,然后将控制权转移给函数。在这种情况下,编译器受到更多限制,因为:
- 它必须将函数to/from使用的变量移动到栈帧
- 它无法在"surrounding"代码的基础上进行其他优化,因为它不知道周围的代码是什么。
__device__
函数必须由编译器以独立方式处理。
因此,如果函数可以内联,那么您的两种方法之间应该没有太大区别。如果函数不能内联,那么上述两种方式在寄存器使用上通常会有明显的区别。
可能影响编译器是否尝试内联 __device__
函数的一些明显因素是:
如果
__device__
函数与调用它的__global__
或其他__device__
函数位于单独的编译单元中。在这种情况下,唯一可行的方法是通过 CUDA separate compilation and linking,也称为设备链接。在这种情况下,编译器将不会(不能)内联函数。如果指定了
__noinline__
compiler directive。请注意,这只是对编译器的提示;它可能会被忽略。