强制CUDA使用寄存器作为变量

我在内核中有很多未使用的寄存器。我想告诉CUDA使用一些寄存器来保存一些数据，而不是每次需要时都读取全局数据。（我不能够使用共享MEM）强制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 
}

挥发性声明创建堆栈存储：
4096字节堆栈帧， 0字节溢出店，二手0字节溢出负荷
21名抵抗者，40个字节CMEM [0]

1）是否有一个简单的方法来告诉编译器使用寄存器空间的变量？
2）'堆栈帧'在哪里：寄存器，全局mem，本地mem，...？什么是堆栈框架？（由于当没有所述GPU具有堆叠的虚拟堆叠？）
3）simple.ptx文件基本上是空的：（NVCC -arch sm_20 -ptx simple.cu）

.loc 2 14 2 
ret;

任何想法，我可以找到真正的机器/编译代码？

来源

2012-08-28 Doug

编译器优化了整个代码，因为它不修改任何非瞬态状态。 – njuffa

每个线程要求1024个寄存器是一个非常高的顺序。大多数内核每个线程需要数十个寄存器。如果你想确保编译器可以使用一个寄存器作为变量，它需要是一个标量（即不是你在'for'循环中索引的数组）。 –

在哪里/什么堆栈框架答案可以在这里找到：http://stackoverflow.com/questions/7810740/where-does-cuda-allocate-the-stack-frame-for-kernels – Doug

动态索引数组无法存储在寄存器中，因为GPU寄存器文件不是动态可寻址的。
标量变量由编译器自动存储在寄存器中。
静态索引（即，其中的索引可以在编译时来确定）小阵列（比如，小于16个浮子）可以被存储在由编译器寄存器。

SM 2.0 GPU（费米）只支持每个线程多达63个寄存器。如果超过此值，寄存器值将从缓存层次结构支持的本地（片外）内存溢出/填充。 SM 3.5 GPU将其扩展到每个线程最多255个寄存器。一般来说，正如Jared所提到的，每个线程使用太多的寄存器是不期望的，因为它减少了占用，因此减少了内核中的延迟隐藏能力。 GPU在并行性方面蓬勃发展，并通过覆盖来自其他线程的工作来延迟内存延迟。

因此，你可能不应该优化阵列到寄存器。相反，请确保跨线程访问这些阵列的内存尽可能接近顺序，以便最大程度地实现合并（即最大限度地减少内存事务）。

你给可以是用于共享存储器如果的情况下的例子：

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

的每线程阵列尺寸足够小，以分配足够的空间用于多个线程块中的所有线程（每个线程1024浮点数远远多）。

正如njuffa提到的那样，你的内核只使用2个寄存器的原因是你没有对内核中的数据做任何有用的事情，死编码全部被编译器清除。

来源

2012-08-29 01:04:01 harrism

您建议线程可以使用的reg的数量是有限制的（SM_20为63）。这是从哪里来的？设备属性显示每个BLOCK（regsPerbBock）的reg的数量限制。 – Doug

它来自体系结构，编译器负责确保没有大于生成的二进制代码中使用的限制的寄存器数量。除了性能方面的原因（例如了解寄存器溢出的原因），用户不必担心此限制，这就是为什么不需要将其列在deviceProps结构中的原因。 – harrism

使用许多寄存器可能是需要的，因为最大化占用并不是隐藏延迟的唯一方法。隐藏延迟的另一种方式是指令级并行。有时它是达到最佳性能的唯一途径。查看瓦西里沃尔科夫[幻灯片]（http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf），其中的作者在只有8％的入住率时获得了高峰表现。 –

如前所述，寄存器（和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用作堆栈指针，因此用于“寄存器溢出”（这意味着通过临时将其值存储在堆栈中释放寄存器，并在更多寄存器需要比可用），所以它是不可接触的。

来源

2012-08-29 20:28:04 Dude

有关-Xopencc = -O3的链接消失了，我无法在CUDA上下文中找到任何引用。你能否指点我一些资源或解释近期cuda（7.0/7.5）的行为是否相似？ – XapaJIaMnu

强制CUDA使用寄存器作为变量

回答

相关问题