15

我的内核中有许多未使用的寄存器。我想告诉 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
}

注册声明什么都不做。
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;

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

4

2 回答 2

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

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

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

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

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

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

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

于 2012-08-29T01:04:01.587 回答
6

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

此外，对于 CUDA 4.0，使用命令行开关-Xopencc=-O3可以将除普通标量（例如数据结构）之外的任何内容映射到寄存器（参见这篇文章）。对于 CUDA > 4.0，您必须禁用调试支持（没有-G命令行选项 - 只有在禁用调试时才会进行优化）。

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

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

于 2012-08-29T20:28:04.227 回答