CUDA：啥时候使用共享内存，啥时候依赖 L1 缓存？

Posted 2023-02-14

【中文标题】CUDA：啥时候使用共享内存，啥时候依赖 L1 缓存？

【英文标题】：CUDA: When to use shared memory and when to rely on L1 caching?CUDA：什么时候使用共享内存，什么时候依赖 L1 缓存？

【发布时间】：2012-07-01 19:25:02

【问题描述】：

Compute Capability 2.0 (Fermi) 发布后，我想知道是否还有任何用于共享内存的用例。也就是说，什么时候使用共享内存比让 L1 在后台发挥它的魔力更好？

共享内存是否只是为了让为 CC

为了通过共享内存进行协作，块中的线程写入共享内存并与__syncthreads()同步。为什么不简单地写入全局内存（通过 L1），并与__threadfence_block() 同步？后一个选项应该更容易实现，因为它不必与两个不同的值位置相关，并且它应该更快，因为没有从全局内存到共享内存的显式复制。由于数据被缓存在 L1 中，线程不必等待数据真正到达全局内存。

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

【参考方案1】：

据我所知，GPU 中的 L1 缓存的行为与 CPU 中的缓存非常相似。因此，您的评论“这与 L1 中的值相反，如果它们不经常使用就会被驱逐”对我来说没有多大意义

L1 缓存上的数据在使用不够频繁时不会被清除。通常，当对以前不在缓存中的内存区域发出请求时，它会被驱逐，并且其地址解析为已在使用的内存区域。我不知道 NVidia 使用的确切缓存算法，但假设一个常规的 n 路关联，那么每个内存条目只能缓存在整个缓存的一小部分中，基于它的地址

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

【讨论】：

谢谢，这确实回答了我的问题。我曾将缓存描绘成能够跟踪哪些元素被使用得最多，并且更喜欢缓存这些元素。我现在已经阅读了 n 路关联缓存，在我看来主要问题是它们可能会抛出一个经常使用的值，因为另一个缓存行适合该插槽。

我认为这意味着编写 CUDA 程序的一个好策略通常可能是首先编写仅使用全局内存的算法，然后查看 L1 是否工作得足够好以隐藏内存延迟。如果算法受内存限制，则考虑使用共享内存进行手动优化。

【参考方案2】：

自动缓存效率低于手动暂存器内存的两大原因（也适用于 CPU）

对随机地址的并行访问效率更高。示例：直方图。假设您要增加 N 个 bin，每个 bin 相隔 > 256 个字节。然后由于合并规则，这将导致 N 次串行读/写，因为全局和高速缓存存储在大约 256 字节的大块中。共享内存没有这个问题。

还要访问全局内存，您必须进行虚拟到物理地址的转换。拥有可以在 || 中进行大量翻译的 TLB会很贵。我还没有看到任何在 || 中实际执行矢量加载/存储的 SIMD 架构。我相信这就是原因。

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

另外，根据NVIDIA employee，当前的 L1 缓存是直写的（立即写入 L2 缓存），这会减慢您的程序。

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

【讨论】：

计算能力 2.* 和 3.* 在写入时使 L1 缓存行无效。计算能力 3.0-3.5 不在 L1 中缓存全局读取。在计算能力 3.* 设备上，每组 8 字节的共享内存带宽实际上是 256 字节/clk，而 L1 被限制为来自高速缓存行的 128 字节。正如耶鲁大学所说，共享内存存在银行冲突（所有访问都必须针对不同的银行或银行中的相同地址），而 L1 存在地址分歧（所有地址必须在同一个 128 字节高速缓存行中），因此共享内存在随机访问。

让我推测一下为什么 SIMD 内存访问在通用处理器上实际上是不存在的（例如，英特尔 AVX2 有一个集合，但它确实是串行的）。我非常确信这是因为进行虚拟到物理地址转换的成本很高，共享内存访问不需要，因为它是它自己的地址空间。想象一下并行执行 32 个 TLB 查找的成本！如果所有 32 个地址都在同一个页面中，也许会有优化？

【参考方案3】：

通过多个内存层缓存数据始终需要遵循缓存一致性协议。有几种这样的协议，而决定哪种协议最合适总是权衡取舍。

你可以看看一些例子：

Related to GPUs

Generally for computing units

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

但是，使线程保持最新，或者使线程保持一致状态需要付出一些代价，这实际上是缺少周期。

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

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

当然，NVIDIA GPU 上的共享内存被分成多个存储区，如果有人想用它来提高性能，应该先看看这个。原因是当两个线程访问同一个银行时会发生银行冲突，这会导致访问序列化...，但这是另一回事link