来自 CUDA Compute Capability 2.0 (Fermi) 的全局内存访问通过 768 KB L2 缓存工作。看起来，开发者不再关心全局内存库了。但是全局内存仍然很慢，所以正确的访问模式很重要。现在的重点是尽可能多地使用/重用 L2。我的问题是，如何？如果我需要一些详细信息，L2 的工作原理以及我应该如何组织和访问全局内存（例如，每个线程 100-200 个元素数组），我将不胜感激。
【问题讨论】：
标签： cuda
L2 缓存在某些方面有所帮助，但它并不能消除对全局内存的合并访问的需要。简而言之，合并访问意味着对于给定的读取（或写入）指令，warp 中的各个线程正在读取（或写入）全局内存中相邻的连续位置，最好是在 128 字节边界上作为一个组对齐.这将最有效地利用可用内存带宽。
在实践中，这通常并不难实现。例如：
int idx=threadIdx.x + (blockDim.x * blockIdx.x);
int mylocal = global_array[idx];
假设global_array 是在全局内存中使用 cudaMalloc 以普通方式分配的，则将在 warp 中的所有线程之间提供合并（读取）访问。这种类型的访问可以 100% 使用可用内存带宽。
一个关键点是内存事务通常发生在 128 字节块中，这恰好是高速缓存行的大小。如果您甚至请求块中的一个字节，则将读取整个块（通常存储在 L2 中）。如果您稍后从该块中读取其他数据，通常会从 L2 对其进行服务，除非它已被其他内存活动逐出。这意味着以下顺序：
int mylocal1 = global_array[0];
int mylocal2 = global_array[1];
int mylocal3 = global_array[31];
通常都由一个 128 字节的块提供服务。 mylocal1 的第一次读取将触发 128 字节读取。对mylocal2 的第二次读取通常会从缓存值（在 L2 或 L1 中）进行服务，而不是通过触发从内存中进行的另一次读取。但是，如果可以适当地修改算法，最好从多个线程连续读取所有数据，如第一个示例所示。这可能只是巧妙地组织数据的问题，例如使用数组结构而不是结构数组。
在许多方面，这类似于 CPU 缓存行为。缓存行的概念以及为来自缓存的请求提供服务的行为都类似。
Fermi L1 和 L2 可以支持回写和直写。 L1 在每个 SM 的基础上可用，并且可配置为与共享内存拆分为 16KB L1（和 48KB SM）或 48KB L1（和 16KB SM）。 L2 跨设备统一，大小为 768KB。
我要提供的一些建议是不要假设 L2 缓存只是修复了草率的内存访问。 GPU 缓存比 CPU 上的等效缓存小得多，因此在那里更容易遇到麻烦。一般的建议是简单地编码，就好像缓存不存在一样。与缓存阻塞等面向 CPU 的策略相比，通常最好将编码工作集中在生成合并访问上，然后在某些特定情况下使用共享内存。然后对于我们无法在所有情况下都进行完美内存访问的不可避免的情况，我们让缓存提供它们的好处。
您可以通过查看一些可用的NVIDIA webinars 来获得更深入的指导。例如，Global Memory Usage & Strategy webinar（和slides）或CUDA Shared Memory & Cache webinar 对本主题具有指导意义。您可能还想阅读CUDA C Programming Guide 的Device Memory Access section。
【讨论】：
__syncthreads();。在上面的例子中，int local1 = global_array[idx]; 后面会跟着__syncthreads();。问题是，如果我有多个像int local1 = global_array1[idx]; int local2 = global_array2[idx]; int local3 = global_array3[idx]; 这样的数组，我会在所有这些定义之后用一个__syncthreads(); 进行合并读取吗？谢谢。
__syncthreads() 没有任何关系。