强制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
}

编译瓦特/:NVCC -arch sm_20 --ptxas选项= -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;

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

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，以便将除普通标量（例如数据结构）以外的任何内容映射到寄存器（请参阅此帖子）。对于 CUDA > 4.0，您必须禁用调试支持（没有-G命令行选项 - 只有在禁用调试时才会进行优化）。

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

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