caching - CUDA:什么时候使用共享内存，什么时候依靠L1缓存？

Question

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

共享内存是否只是存在于其中，以使为CC <2.0设计的算法无需修改即可高效运行？

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

使用共享内存，可以确保在块的整个过程中都将在那里存放的值保留在那里。这与L1中的值相反，L1中的值如果使用不频繁会被逐出。在任何情况下，最好还是让这种很少使用的数据缓存在共享内存中，而不是让L1根据该算法实际使用的模式来管理它们？

Answer 1

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

如果不经常使用L1缓存中的数据，则不会将其驱逐出去。通常，当请求一个以前不在高速缓存中且其地址解析为已在使用中的内存区域的请求时，将其逐出。我不知道NVidia使用的确切的缓存算法，但是假设使用常规的n向关联，则每个内存条目只能基于其地址缓存在整个缓存的一小部分中

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