Compute Capability 2.0 (Fermi) 发布后，我想知道是否还有共享内存的用例。也就是说，什么时候使用共享内存比让 L1 在后台执行它的魔力更好？
共享内存是否只是为了让为 CC < 2.0 设计的算法无需修改即可高效运行？
为了通过共享内存进行协作，块中的线程写入共享内存并与__syncthreads(). 为什么不简单地写入全局内存（通过 L1），并与 同步__threadfence_block()？后一个选项应该更容易实现，因为它不必与值的两个不同位置相关，并且它应该更快，因为没有从全局内存到共享内存的显式复制。由于数据被缓存在 L1 中，线程不必等待数据真正到达全局内存。
使用共享内存，可以保证放置在那里的值在整个块的持续时间内保持在那里。这与 L1 中的值相反，如果不经常使用它们就会被驱逐。在任何情况下，将这些很少使用的数据缓存在共享内存中比让 L1 根据算法实际具有的使用模式来管理它们更好吗？
自动缓存效率低于手动暂存器内存的两大原因（也适用于 CPU）
同样要访问全局内存，您必须进行虚拟到物理地址的转换。拥有可以在 || 中进行大量翻译的 TLB 会很贵。我还没有看到任何在 || 中实际执行矢量加载/存储的 SIMD 架构。我相信这就是原因。
此外，根据NVIDIA 员工的说法，当前的 L1 缓存是直写的（立即写入 L2 缓存），这会降低您的程序速度。
所以基本上，如果你真的想要性能，缓存会妨碍你。
据我所知，GPU 中的 L1 缓存的行为与 CPU 中的缓存非常相似。因此，您的评论“这与 L1 中的值相反，如果它们不经常使用就会被驱逐”对我来说没有多大意义
L1 缓存上的数据在使用不够频繁时不会被驱逐。通常，当对以前不在缓存中的内存区域发出请求时，它会被驱逐，并且其地址解析为已在使用的内存区域。我不知道 NVidia 使用的确切缓存算法，但假设一个常规的 n 路关联，那么每个内存条目只能缓存在整个缓存的一小部分中，基于它的地址
我想这也可以回答你的问题。使用共享内存，您可以完全控制存储内容的位置，而使用缓存，一切都会自动完成。尽管编译器和 GPU 在优化内存访问方面仍然非常聪明，但有时您仍然可以找到更好的方法，因为您知道将给出什么输入，以及哪些线程将执行什么操作（在一定程度上当然程度）
通过多个内存层缓存数据始终需要遵循缓存一致性协议。有几种这样的协议，决定哪一个最合适总是权衡取舍。
你可以看看一些例子：
我不想详细说明，因为这是一个巨大的领域，我不是专家。我想指出的是，在共享内存系统中（这里的术语共享不是指所谓的 GPU 共享内存），其中许多计算单元（CU）同时需要数据，有一个内存协议试图将数据保持在单元附近，以便尽可能快地获取它们。在 GPU 的示例中，当同一个 SM（对称多处理器）中的许多线程访问相同的数据时，应该存在一致性，即如果线程 1 从全局内存中读取一个字节块，并且在下一个周期中线程 2 是如果要访问这些数据，那么一个有效的实现将是线程 2 知道数据已经在 L1 缓存中找到并且可以快速访问它。这就是缓存一致性协议试图实现的目标，让所有计算单元与缓存 L1、L2 等中存在的数据保持同步。
但是，使线程保持最新，或者使线程保持一致状态，需要付出一些代价，这基本上是缺少周期。
在 CUDA 中，通过将内存定义为共享而不是 L1 缓存，您可以将其从一致性协议中释放出来。因此，访问该内存（物理上与任何材料相同）是直接的，并且不会隐式调用一致性协议的功能。
我不知道这应该有多快，我没有执行任何此类基准测试，但我的想法是，由于您不再为此协议付费，因此访问应该更快！
当然，英伟达 GPU 上的共享内存被拆分为银行，如果有人想用它来提高性能，应该先看看这个。原因是当两个线程访问同一个银行时发生银行冲突，这会导致访问序列化......，但这是另一回事 链接