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

注册 声明什么都不做。
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 在并行性上茁壮成长，并通过用其他线程的工作覆盖内存延迟来实现这一点。

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

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

块中的许多线程使用相同的数据，或

每个线程的数组大小足够小，可以为多个线程块中的所有线程分配足够的空间(每个线程 1024 个浮点数远远不够)。

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