CUDA:合并全局内存访问速度比共享内存快吗？另外,分配大型共享内存阵列会减慢程序的速度吗？

Question

我没有发现NVIDIA Tesla M2050的共享内存速度有所改善,每块大约有49K共享内存.实际上,如果我在共享内存中分配一个大的char数组,它会减慢我的程序.例如

__shared__ char database[49000];

让我比较慢的运行时间

__shared__ char database[4900];

程序只访问数据库的前100个字符,因此不需要额外的空间.我无法弄清楚为什么会这样.任何帮助,将不胜感激.谢谢.

Answer 1

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

每个多处理器托管多个处理器; 对于现代设备,通常为32,经线中的线程数.这意味着,在没有分歧或存储器停顿的情况下,平均处理速率是每个周期32个指令(由于流水线操作,延迟很高).

CUDA为多处理器安排了几个块.每个块由几个经线组成.当warp在全局内存访问上停止(即使合并的访问具有高延迟),也会处理其他warp.这有效地隐藏了延迟,这就是GPU中可接受高延迟全局内存的原因.为了有效地隐藏延迟,您需要执行足够的额外warp,直到停止的warp可以继续.如果所有warp在内存访问中停止,则无法再隐藏延迟.

共享内存分配给CUDA中的块,并存储在GPU设备上的单个多处理器中.每个多处理器都有一个相对较小的固定数量的共享内存空间.与多处理器在共享内存和寄存器使用方面可支持的相比,CUDA无法为多处理器安排更多块.换句话说,如果多处理器上的共享内存量是X并且每个块都需要Y共享内存,那么CUDA将一次安排不超过底层(X/Y)块到每个多处理器(它可能会更少,因为有其他约束,例如寄存器使用).

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

编辑:

为了解决问题的其他部分,假设没有共享内存库冲突和完全合并全局内存访问......这个答案有两个方面:延迟和带宽.共享内存的延迟将低于全局内存的延迟,因为共享内存在片上.带宽将大致相同.因此,如果您能够通过合并隐藏全局内存访问延迟,则不会受到惩罚(注意:访问模式在这里非常重要,因为共享内存允许可能更多样化的访问模式,几乎没有性能损失,因此可以即使您可以隐藏所有全局内存延迟,也可以使用共享内存.