CUDA:合并的全局内存访问是否比共享内存快？另外，分配大的共享内存阵列是否会降低程序速度？

Question

我在NVIDIA Tesla M2050上发现共享内存的速度没有提高
每个块约有49K共享内存。其实如果我分配
共享内存中的大字符数组会减慢我的程序的速度。例如

__shared__ char database[49000];

给我的运行时间比

__shared__ char database[4900];

该程序仅访问数据库的前100个字符，因此额外的空间
是没有必要的。我不知道为什么会这样。任何帮助，将不胜感激。
谢谢。

Answer 1

当使用较大的阵列时，CUDA共享内存性能相对较差的原因可能与以下事实有关:每个多处理器具有有限数量的可用共享内存。

每个多处理器托管几个处理器；对于现代设备(通常为32个)，扭曲中的线程数。这意味着，在没有差异或内存停顿的情况下，平均处理速率为每个周期32条指令(由于流水线而导致等待时间较长)。

CUDA将几个块调度到一个多处理器。每个块由几个变形组成。当warp在全局内存访问上停顿时(即使合并的访问也具有高延迟)，其他warp也会被处理。这有效地隐藏了延迟，这就是为什么在GPU中可以接受高延迟全局内存的原因。为了有效地隐藏延迟，您需要执行足够的额外扭曲，直到停滞的扭曲可以继续。如果所有扭曲都因内存访问而停滞，则您将无法再隐藏延迟。

共享内存分配给CUDA中的块，并存储在GPU设备上的单个多处理器上。每个多处理器具有相对较小的固定数量的共享内存空间。就共享内存和寄存器使用而言，CUDA不能为多处理器调度比多处理器可以支持的块更多的块。换句话说，如果多处理器上的共享内存量为X，并且每个块需要Y共享内存，则CUDA一次向每个多处理器调度的层数不超过floor(X/Y)个块(这可能会少一些，因为其他限制，例如寄存器使用情况)。

因此，通过增加块的共享内存使用率，您可能会减少内核的事件线程数(占用率)，从而损害性能。您应该使用-Xptxas =“-v”标志进行编译，以查看内核代码。这应该为您提供每个内核的注册，共享和恒定内存使用率。在最新版本的CUDA占用率计算器中，使用此数据和您的内核启动参数以及其他必需的信息来确定您是否可能受到占用率的影响。

编辑:

为了解决您的问题的另一部分，假设没有共享内存库冲突，并且全局内存访问完美合并，那么这个答案有两个方面:延迟和带宽。由于共享内存在芯片上，因此共享内存的延迟将低于全局内存的延迟。带宽将几乎相同。因此，如果您能够通过合并隐藏全局内存访问延迟，则不会有任何损失(请注意:访问模式在这里很重要，因为共享内存允许潜在的更多样化的访问模式，而几乎没有性能损失，因此可以即使您可以隐藏所有全局内存延迟，也可以使用共享内存。