为什么 OpenCL Nvidia 编译器 (nvcc) 不使用寄存器两次？

Question

我正在使用 Nvidia 驱动程序做一个小型 OpenCL 基准测试，我的内核执行 1024 次保险丝乘加并将结果存储在一个数组中：

#define FLOPS_MACRO_1(x)    { (x) = (x) * 0.99f + 10.f; } // Multiply-add
#define FLOPS_MACRO_2(x)    { FLOPS_MACRO_1(x) FLOPS_MACRO_1(x) }
#define FLOPS_MACRO_4(x)    { FLOPS_MACRO_2(x) FLOPS_MACRO_2(x) }
#define FLOPS_MACRO_8(x)    { FLOPS_MACRO_4(x) FLOPS_MACRO_4(x) }
// more recursive macros ...
#define FLOPS_MACRO_1024(x) { FLOPS_MACRO_512(x) FLOPS_MACRO_512(x) }

__kernel void ocl_Kernel_FLOPS(int iNbElts, __global float *pf)
{
   for (unsigned i = get_global_id(0); i < iNbElts; i += get_global_size(0))  
   {
      float f = (float) i;
      FLOPS_MACRO_1024(f)
      pf[i] = f;
    }   
}

但是当我查看生成的 PTX 时，我看到了：

    .entry ocl_Kernel_FLOPS(
    .param .u32 ocl_Kernel_FLOPS_param_0,
    .param .u32 .ptr .global .align 4 ocl_Kernel_FLOPS_param_1
)
{
    .reg .f32   %f<1026>; // 1026 float registers !
    .reg .pred  %p<3>;
    .reg .s32   %r<19>;    

    ld.param.u32    %r1, [ocl_Kernel_FLOPS_param_0];
    // some more code unrelated to the problem
    // ...

BB1_1:
    and.b32     %r13, %r18, 65535;
    cvt.rn.f32.u32  %f1, %r13;
    fma.rn.f32  %f2, %f1, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f3, %f2, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f4, %f3, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f5, %f4, 0f3F7D70A4, 0f41200000;
    // etc
    // ...

如果我是对的，PTX 使用1026 个浮点寄存器来执行 1024 个操作，并且即使它可以仅使用 2 个寄存器执行所有乘加运算，也永远不会重用寄存器两次。1026 远高于线程允许拥有的最大寄存器数（根据规范），所以我猜这最终会导致内存溢出。

这是编译器错误还是我完全错过了什么？

我在 Quadro K2000 GPU 上使用 nvcc 6.5 版。

编辑

实际上我确实错过了规范中的一些内容：

“由于 PTX 支持虚拟寄存器，编译器前端生成大量寄存器名称是很常见的。而不是要求显式声明每个名称，PTX 支持创建一组变量的语法，这些变量具有附加的公共前缀字符串整数后缀。例如，假设一个程序使用了大量的 .b32 变量，比如一百个，命名为 %r0, %r1, ..., %r99"

Answer 1

的PTX文件格式是用来描述一个虚拟机和指令集架构：

PTX 定义了用于通用并行线程执行的虚拟机和 ISA。PTX 程序在安装时被翻译成目标硬件指令集。PTX 到 GPU 转换器和驱动程序使 NVIDIA GPU 能够用作可编程并行计算机。

因此，您在那里获得的 PTX 输出不是“GPU 汇编程序”的形式。它只是一种中间表示，旨在能够描述几乎任何形式的并行计算。

然后将 PTX 表示编译为相应目标 GPU 的实际二进制文件。为了能够从实际架构中抽象出来，这很重要- 特别是对于您的示例：应该可以使用程序的相同PTX 表示，而不管特定目标机器上可用的寄存器数量如何。您看到的 1026 个“寄存器”是“虚拟”寄存器，最终可能会映射到实际可用的（少数）真实硬件寄存器。您可以--ptxas-options=-v在编译期间将参数添加到 NVCC 以获取有关寄存器使用的附加信息。

（这与LLVM背后的想法大致相同——即，拥有一个可以优化和争论的表示，既从原始源代码中抽象出来，也从实际目标架构中抽象出来）。