优化点积中的寄存器使用

Question

我正在开发一个具有多个向量运算（如标量和向量积）的核函数。内核使用了大量的寄存器，因此占用率很低。我正在尝试减少使用的寄存器数量以提高占用率。

考虑以下 __device__ 函数在两个 float3 之间执行标量积：

__device__ float dot(float3 in1, float3 in2) { return in1.x * in2.x + in1.y * in2.y + in1.z * in2.z; }

如果我使用

生成 .ptx 文件

nvcc -ptx -gencode arch=compute_52,code=sm_52 -rdc=true simpleDot2.cu

（文件simpleDot2.cu只包含__device__函数的定义），我基本上得到

    // .globl   _Z3dot6float3S_
.visible .func  (.param .b32 func_retval0) _Z3dot6float3S_(
    .param .align 4 .b8 _Z3dot6float3S__param_0[12],
    .param .align 4 .b8 _Z3dot6float3S__param_1[12]
)
{
    .reg .f32   %f<10>;


    ld.param.f32    %f1, [_Z3dot6float3S__param_0+8];
    ld.param.f32    %f2, [_Z3dot6float3S__param_0];
    ld.param.f32    %f3, [_Z3dot6float3S__param_0+4];
    ld.param.f32    %f4, [_Z3dot6float3S__param_1+8];
    ld.param.f32    %f5, [_Z3dot6float3S__param_1];
    ld.param.f32    %f6, [_Z3dot6float3S__param_1+4];
    mul.f32     %f7, %f3, %f6;
    fma.rn.f32  %f8, %f2, %f5, %f7;
    fma.rn.f32  %f9, %f1, %f4, %f8;
    st.param.f32    [func_retval0+0], %f9;
    ret;
}

从.ptx代码看，似乎用到了9个寄存器，或许可以降低。我明白 .ptx 代码不是 GPU 执行的最终代码。

问题

是否有机会重新安排.ptx代码中的寄存器使用，例如回收寄存器f1-f6，以减少整体数量占用的寄存器数？

非常感谢您的帮助。

Answer 1

TL;DR 第一次订购，没有。

PTX既是虚拟ISA又是编译器中间表示。 PTX代码中使用的寄存器是虚拟寄存器，与GPU的物理寄存器没有固定关系。 CUDA 工具链生成的 PTX 代码遵循 SSA（静态单一赋值）约定。这意味着每个虚拟寄存器只被写入一次。换句话说：当一条指令产生一个结果时，它被分配给一个新的寄存器。这意味着更长的内核可能会使用数千个寄存器。

在 CUDA 工具链中，PTX 代码由 ptxas 组件编译为机器代码 (SASS)。因此，尽管名称如此，但这并不是一个汇编程序，而是一个 optimizing 编译器，可以执行循环展开、CSE（公共子表达式消除）等操作。最重要的是，ptxas 负责寄存器分配和指令调度，以及针对特定 GPU 架构的所有优化。

因此，任何对寄存器使用问题的检查都需要关注机器代码，可以使用 cuobjdump --dump-sass 提取机器代码。此外，程序员对使用的寄存器数量的影响非常有限，因为 ptxas 在确定寄存器分配时使用大量试探法，特别是在寄存器使用与性能之间进行权衡：提前调度加载往往会增加寄存器压力寿命范围，在 CSE 期间创建临时变量或创建归纳变量以降低循环强度。

以 3.0 及更高计算能力为目标的现代 CUDA 版本在确定这些 trade-offs 时通常会做出很好的选择，程序员很少需要考虑寄存器压力。目前尚不清楚是什么激发了提问者在这方面的问题。

CUDA 中记录的控制最大寄存器使用的机制是 -maxrregcount command-line 标志 nvcc，它适用于整个编译单元，以及 __launch_bounds__允许在 per-kernel 基础上进行控制的属性。有关详细信息，请参阅 CUDA 文档。除此之外，可以尝试通过选择 pxtas 优化级别和 -Xptxas -O{1|2|3}（默认为 -O3）或 re-arranging 源代码，或使用倾向于简化生成代码的编译器标志，例如 -use_fast_math.

当然，这种间接方法可能会产生许多其他通常无法预测的效果，任何理想的结果都将是 "brittle"，例如通过更改为新版本的工具链很容易销毁。

优化点积中的寄存器使用

Optimizing register usage in dot product

cuda

ptx