为什么 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 defines a virtual machine and ISA for general purpose parallel thread execution. PTX programs are translated at install time to the target hardware instruction set. The PTX-to-GPU translator and driver enable NVIDIA GPUs to be used as programmable parallel computers.

所以您在那里获得的 PTX 输出是不是 "GPU assembler" 的一种形式。它只是一种中间表示，旨在能够描述几乎任何形式的并行计算。

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

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

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

Why doesn't OpenCL Nvidia compiler (nvcc) use the registers twice?

nvidia

opencl

nvcc

ptx