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

Question

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

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

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

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

Answer 1

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

如果不经常使用L1高速缓存的数据，则不会将其逐出。通常在对先前不在高速缓存中的内存区域发出请求并且其地址解析为已在使用的内存区域时被逐出。我不知道NVidia使用的确切缓存算法，但假设有一个常规的n路关联，那么每个内存条目只能根据它的地址缓存在整个缓存的一小部分中。

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

Answer 2

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

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

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

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

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

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