CUDA:何时使用共享内存以及何时依赖L1缓存？

Question

CUDA:何时使用共享内存以及何时依赖L1缓存？

在Compute Capability 2.0(Fermi)发布之后,我想知道是否有任何用例共享内存.也就是说,何时使用共享内存比让L1在后台执行魔术更好？

共享内存是否只是为了让CC <2.0设计的算法在没有修改的情况下高效运行？

要通过共享内存进行协作,块中的线程会写入共享内存并与之同步__syncthreads().为什么不简单地写入全局内存(通过L1),并与__threadfence_block()？同步？后一个选项应该更容易实现,因为它不必涉及值的两个不同位置,并且它应该更快,因为没有从全局到共享内存的显式复制.由于数据在L1中缓存,因此线程不必等待数据实际上一直到全局内存.

使用共享内存,可以保证在整个块的持续时间内放置的值保持不变.这与L1中的值相反,如果它们不经常使用,则会被逐出.是否有任何情况下,在共享内存中缓存这些很少使用的数据比让L1根据算法实际具有的使用模式管理它们更好？

Answer 1

Yal*_*ang 10

自动缓存效率低于手动便笺本内存的两大原因(也适用于CPU)

对随机地址的并行访问更有效.示例:直方图.假设你想增加N个bin,每个bin的间隔大于256个字节.然后由于合并规则,这将导致N个串行读/写,因为全局和高速缓存存储器被组织成大约256字节的块.共享内存没有这个问题.

另外,要访问全局内存,您必须进行虚拟到物理地址转换.拥有一个可以在||中进行大量翻译的TLB 将是相当昂贵的.我还没有看到任何实际上在||中进行矢量加载/存储的SIMD架构我相信这就是原因.

避免将死值写回内存,这会浪费带宽和功率.示例:在图像处理管道中,您不希望将中间图像刷新到内存中.

此外,根据NVIDIA员工的说法,当前的L1缓存是直写(立即写入L2缓存),这将减慢您的程序.

所以基本上,如果你真的想要性能,缓存就会受到阻碍.

计算能力 2.* 和 3.* 在写入时使 L1 缓存行无效。计算能力 3.0-3.5 不在 L1 中缓存全局读取。在计算能力 3.* 设备上，每组 8 字节的共享内存带宽实际上是 256 字节/时钟，而 L1 限制为来自缓存线的 128 字节。正如 Yale 所说，共享内存有 bank 冲突（所有访问必须是对不同 bank 或 bank 中的相同地址）而 L1 有地址分歧（所有地址必须在同一个 128 字节缓存行中）所以共享内存在随机访问。 (2认同)

Answer 2

Nap*_*s62 8

据我所知,GPU中的L1缓存表现得与CPU中的缓存非常相似.因此,您的评论"这与L1中的值相反,如果它们不经常使用就被驱逐"对我来说没有多大意义

如果不经常使用L1高速缓存的数据,则不会将其逐出.通常在对先前不在高速缓存中的内存区域发出请求并且其地址解析为已在使用的内存区域时被逐出.我不知道由NVIDIA所采用的确切缓存算法,但假设一个普通N路相连,那么每个存储条目只能在整个高速缓存的一小部分缓存,基于它的地址

我想这也可以回答你的问题.使用共享内存,您可以完全控制存储在哪里,而对于缓存,一切都是自动完成的.即使编译器和GPU在优化内存访问方面仍然非常聪明,你有时仍然可以找到更好的方法,因为你是知道将给出什么输入的人,以及什么线程会做什么(对某些当然程度)

我认为这意味着编写CUDA程序的好策略通常可能是首先编写算法以仅使用全局内存,并查看L1是否运行良好以至于隐藏了内存延迟.如果算法结果是内存绑定,则考虑使用共享内存进行手动优化. (3认同)

归档时间：	13 年，8 月前
查看次数：	6832 次
最近记录：	7 年，5 月前