为什么 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 远远高于线程允许拥有的最大寄存器数（根据specs），所以我猜这最终会导致内存溢出。

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

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

编辑

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

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

Answer 1

PTX file format 旨在描述虚拟机和指令集架构：

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

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

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

（这与LLVM 背后的想法大致相同 - 即，有一个可以优化和争论的表示，两者都从原始源代码抽象和实际目标架构）。