gpt4 book ai didi

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

转载 作者:行者123 更新时间:2023-12-02 04:15:31 24 4
gpt4 key购买 nike

我在很多地方读到过,__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 。如果它是内联的,这通常会给优化编译器(主要是 ptxas )最好的机会来调整寄存器的使用,因为它认为合适。 (请注意,至少在理论上,这种“适应”可能会导致使用更多或更少的寄存器。但是,内联情况通常会导致编译器使用更少的寄存器并可能提高性能。但是编译器主要针对更高的性能进行优化,而不是更少的寄存器使用。)

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

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

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

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

  1. 如果 __device__ 函数位于与调用它的 __global__ 或其他 __device__ 函数不同的编译单元中。在这种情况下,唯一可行的方法是通过 CUDA separate compilation and linking ,也称为设备链接。在这种情况下,编译器不会(无法)内联该函数。

  2. 如果 __noinline__ compiler directive已指定。请注意,这只是对编译器的提示;它可能会被忽略。

关于performance - 调用 __device__ 函数会影响 CUDA 中使用的寄存器数量吗?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/34046227/

24 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com