强制 CUDA 将寄存器用于变量

Question

我的内核中有许多未使用的寄存器。我想告诉 CUDA 使用一些寄存器来保存一些数据，而不是每次需要时都读取全局数据。 （我无法使用共享内存。）

__global__ void simple(float *gData) {
float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

编译 w/: nvcc -arch sm_20 --ptxas-options=-v simple.cu，我得到
0 字节堆栈帧，0 字节溢出存储，0 字节溢出加载
使用了 2 个寄存器，40 字节 cmem[0]

__global__ void simple(float *gData) {
register float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

register 声明什么都不做。
0 字节堆栈帧，0 字节溢出存储，0 字节溢出加载
使用了 2 个寄存器，40 字节 cmem[0]

__global__ void simple(float *gData) {
volatile float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

volatile 声明创建堆栈存储：
4096 字节堆栈帧，0 字节溢出存储，0 字节溢出加载
使用了 21 个电阻，40 字节 cmem[0]

1) 有没有一种简单的方法可以告诉编译器为变量使用寄存器空间？
2）“堆栈框架”在哪里：寄存器，全局内存，本地内存，...？什么是栈帧？ （GPU 什么时候有堆栈？虚拟堆栈？）
3）simple.ptx文件基本为空：(nvcc -arch sm_20 -ptx simple.cu)

.loc 2 14 2
ret;

知道在哪里可以找到真机/编译代码吗？

Answer 1

• 动态索引数组无法存储在寄存器中，因为 GPU 寄存器文件不可动态寻址。
• 编译器会自动将标量变量存储在寄存器中。
• 静态索引（即可以在编译时确定索引的位置），小型数组（例如，小于 16 个浮点数）可能 由编译器存储在寄存器中。

SM 2.0 GPU (Fermi) 仅支持每个线程最多 63 个寄存器。如果超过这个值，寄存器值将从本地（片外）内存溢出/填​​充，由缓存层次结构支持。 SM 3.5 GPU 将其扩展到每个线程最多 255 个寄存器。

一般来说，正如 Jared 所提到的，每个线程使用过多的寄存器是不可取的，因为它会降低占用率，从而降低内核中的延迟隐藏能力。 GPU 在并行性上蓬勃发展，并通过使用其他线程的工作来覆盖内存延迟来实现这一点。

因此，您可能不应该将数组优化为寄存器。相反，请确保跨线程对这些数组的内存访问尽可能接近顺序，以便最大化合并（即最小化内存事务）。

您给出的示例可能是共享内存的情况如果：

1. 块中的许多线程使用相同的数据，或者
2. 每个线程的数组大小足够小，可以为多个线程块中的所有线程分配足够的空间（每个线程 1024 个浮点数非常多）。

正如 njuffa 提到的，你的内核只使用 2 个寄存器的原因是因为你没有对内核中的数据做任何有用的事情，并且死代码都被编译器消除了。

Answer 2

如前所述，寄存器（和 PTX“参数空间”）不能动态索引。为了做到这一点，编译器必须发出 switch...case 块的代码，以将动态索引转换为立即数。我不确定它是否会自动执行。您可以使用固定大小的元组结构和switch...case 来帮助它发生。 C/C++ 元编程很可能是使此类代码易于管理的首选武器。

此外，对于 CUDA 4.0，使用命令行开关 -Xopencc=-O3 以便将除普通标量（例如数据结构）之外的任何内容映射到寄存器（请参阅 this post）。对于 CUDA > 4.0，您必须禁用调试支持（无 -G 命令行选项 - 仅在禁用调试时进行优化）。

PTX 级别允许比硬件更多的虚拟 寄存器。这些在加载时映射到硬件寄存器。您指定的寄存器限制允许您设置生成的二进制文件使用的硬件资源的上限。当编译到 PTX 时，它可以作为编译器决定何时溢出（见下文）寄存器的启发式方法，以便可以满足某些并发需求（参见 CUDA 文档中的“启动边界”、“占用”和“并发内核执行” - 你可能也会喜欢this most interesting presentation)。

对于 Fermi GPU，最多有 64 个硬件寄存器。第 64 个（或最后一个 - 当使用小于硬件的最大值时）被 ABI 用作堆栈指针，因此用于“寄存器溢出”（这意味着通过将寄存器的值临时存储在堆栈上来释放寄存器，并在更多寄存器时发生需要而不是可用），所以它是不可触碰的。