CUDA：何时使用共享内存，何时依赖于L1缓存？

Question

在Compute Capability 2.0（Fermi）发布之后，我一直在想是否还有任何用例可以使用共享内存。也就是说，在什么情况下使用共享内存比让L1在后台执行更好？

共享内存只是为了让设计针对CC < 2.0的算法能够高效运行而存在吗？

要通过共享内存进行协作，块中的线程将写入共享内存并与__syncthreads()同步。为什么不直接写入全局内存（通过L1），并使用 __threadfence_block()进行同步呢？后者应该更容易实现，因为它不必涉及到两个不同值的位置，并且应该更快，因为没有从全局内存显式复制到共享内存。由于数据被缓存在L1中，线程不必等待数据实际上全部传输到全局内存。

使用共享内存，可以保证放置在其中的值在整个块的持续时间内保持在那里。这与L1中的值相反，如果它们不经常使用，则会被清除。是否有任何情况需要将这些很少使用的数据缓存在共享内存中，而不是根据算法实际使用模式让L1管理它们？

- Roger Dahl

Answer 1

自动缓存比手动刮板内存（适用于CPU）效率低的两个主要原因：

1. 并行访问随机地址更有效。例如：直方图。假设您想增加N个bin，每个bin相距>256字节。由于合并规则，这将导致N个串行读/写，因为全局和缓存内存是以大约256字节的块组织的。共享内存没有这个问题。

   此外，要访问全局内存，必须进行虚拟到物理地址转换。拥有可以同时执行多个转换的TLB将非常昂贵。我还没有看到任何实际上在||中执行向量加载/存储的SIMD架构，我认为这就是原因。

2. 它避免将死值写回内存，从而浪费带宽和功率。例如：在图像处理管道中，您不希望将中间图像刷新到内存。

   此外，根据NVIDIA员工的说法，当前的L1缓存是写通的（立即写入L2缓存），这会减慢程序的速度。

基本上，如果你真的想要性能，缓存会妨碍你。

- Yale Zhang

---

计算能力为2.*和3.*的设备在写操作时会使L1高速缓存行无效。计算能力为3.0-3.5的设备不会将全局读取缓存到L1中。在计算能力为3.*的设备上，每个银行具有8字节共享内存带宽，实际上是256字节/时钟周期，而L1仅限于来自缓存行的128字节。正如耶鲁大学所述，共享内存具有银行冲突（所有访问必须访问不同的银行或同一银行中的相同地址），而L1具有地址分歧（所有地址必须位于同一128字节的缓存行中），因此共享内存在随机访问方面更加有效率。 - Greg Smith

让我就为什么通用处理器上的SIMD内存访问实际上不存在提出一个猜测（例如，Intel AVX2有一个gather，但它实际上是串行的）。我相信这是因为进行虚拟到物理地址转换的成本很高，而共享内存访问不需要这样做，因为它具有自己的地址空间。想象一下同时执行32个TLB查找的成本！也许如果所有32个地址都在同一页中，就会有优化？ - Yale Zhang

Answer 2

据我所知，GPU中的L1缓存与CPU中的缓存非常相似。因此，你的评论“这与L1中的值相反，如果它们没有被足够频繁地使用，则会被清除”对我来说没有多少意义。

当L1缓存中没有使用某个内存区域时，不会将其清除。通常情况下，当请求一个之前未在缓存中的内存区域，并且其地址已经解析为已经在使用的地址时，才会将其清除。我不知道NVidia采用了什么样的缓存算法，但假设采用了常规的n路关联，则每个内存条目只能被缓存在整个缓存的一小部分中，基于它的地址。

我想这也可能回答了你的问题。使用共享内存时，您可以完全控制何处存储数据，而使用缓存时，所有操作都是自动完成的。尽管编译器和GPU仍然可以非常聪明地优化内存访问，但有时您仍然可以找到更好的方法，因为您知道将提供什么输入，以及哪些线程将执行什么操作（当然是在一定程度上）。

- Naps62

---

谢谢，这回答了我的问题。我原本认为缓存可以跟踪哪些元素被最多使用，并且更喜欢缓存它们。现在我已经阅读了关于n路关联缓存的资料，看起来主要问题是它们可能会丢弃一个经常使用的值，只是因为另一条缓存线适合那个插槽。 - Roger Dahl

我认为这意味着编写CUDA程序的好策略通常是先编写只使用全局内存的算法，看看L1缓存是否足够好以隐藏内存延迟。然后，如果算法被限制于内存，再考虑使用共享内存进行手动优化。 - Roger Dahl

Answer 3

通过几层存储器缓存数据时，总是需要遵循一种缓存一致性协议。有几种这样的协议，决定哪一种最适合总是一个权衡。

您可以查看一些示例：

• 与GPU相关：GPU架构的Cache Coherence（DOI：10.1109/HPCA.2013.6522351）

• 通常用于计算单元：芯片上缓存一致性协议概述（DOI：10.1109/IntelliSys.2017.8324309）

我不想深入细节，因为这是一个庞大的领域，而且我也不是专家。我想指出的是，在共享内存系统中（这里的“共享”一词并不是指GPU的所谓共享内存），当许多计算单元（CUs）需要同时访问数据时，存在一种内存协议，试图将数据保持靠近单元，以便尽可能快地获取它们。在GPU的例子中，当同一SM（对称多处理器）中的许多线程访问相同的数据时，应该有一致性，即如果线程1从全局内存读取一块字节，并且在下一个周期线程2将要访问这些数据，则有效的实现应该是线程2知道数据已经在L1缓存中找到，并且可以快速访问它。这就是缓存一致性协议试图实现的，让所有计算单元了解缓存L1、L2等中存在哪些数据。

然而，保持线程最新状态，或者说保持线程处于一致状态，会带来一些成本，这本质上是缺失的周期。

在CUDA中，通过将内存定义为共享内存而不是L1缓存，您可以使其摆脱一致性协议。因此，对该内存的访问（物理上是相同的任何材料）是直接的，并且不会隐式调用一致性协议的功能。

我不知道这应该有多快，我没有进行任何基准测试，但是想法是，由于您不再为此协议付费，因此访问速度应该更快！

当然，NVIDIA GPU上的共享内存被分成了几个bank，如果有人想要使用它来提高性能，应该在此之前查看一下。原因是银行冲突，当两个线程访问同一个bank时，这会导致访问序列化...，但那是另一回事。link

- Andreas Hadjigeorgiou