nvcc 是否优化寄存器使用？

Question

我有以下内核：

    void version1(float *X, float *Y, int N) {
        int n;
        float x,y;

        n = blockIdx.x * blockDim.x + threadIdx.x;
        if (n >= N) return;

        x=X[n];
       x=x+1;
       X[n]=x;

       y=Y[n];
       y=y+1;
       Y[n]=y;
    }

和第二个版本

    void version2(float *X, float *Y, int N) {
        int n;
        float Xb47w;

        n = blockIdx.x * blockDim.x + threadIdx.x;
        if(n >= N) return;

        Xb47w=X[n];
        Xb47w=Xb47w+1;
        X[n]=Xb47w;

        Xb47w=Y[n];
        Xb47w=Xb47w+1;
        Y[n]=Xb47w;
    }

它们产生相同的结果。然而，版本 1 更容易阅读，而版本 2 更难阅读，因为 Xb47w 用于 X 以及 Y。所以我更喜欢版本 1，但是版本 2 有两个寄存器 x y 而不是 1 Xb47w。我有很多内核以这种方式保存寄存器，但更难阅读和维护。

x 在 X[n]=x 之后不再使用所以我想知道 CUDA 编译器是否理解这一点并使版本 1 与版本 2 几乎相同，从而节省一个寄存器？

Answer 1

在内部，nvcc 使用 C++ 编译器来优化代码（嗯，我过于简单化了 link）因此，问题是 C++ 编译器会保存寄存器吗？

答案是使用 godbolt 并比较两个程序的汇编！

编辑：这不是全部，您将看到的是程序的 PTX 表示（也可以使用 nvcc 获得）。下一步是查看称为 SASS 的 gpu 组件本身（取决于卡）。

Answer 2

Does nvcc optimize register usage?

是的，nvcc 试图编译您的代码以使用更少的寄存器（尽管最小寄存器使用本身并不是目标）。

I wonder if the CUDA compiler understands that and makes version1 nearly identical to version2, thus saving one register?

是的，确实如此。或者更确切地说，它不“理解”您的代码的作用，但它注意到冗余 variables/values 并在优化过程中将其删除。

因此，函数的两个版本 compile to the same PTX code (GodBolt.org)

nvcc 是否优化寄存器使用？

Does nvcc optimize register usage?

cuda

compiler-optimization