2

我正在使用 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"

4

1 回答 1

5

PTX 文件格式旨在描述虚拟机和指令集架构：

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

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

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

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

于 2015-02-27T14:47:47.253 回答