CUDA 计算能力 2.0。全局内存访问模式

Question

来自 CUDA Compute Capability 2.0 (Fermi) 的全局内存访问通过 768 KB L2 缓存工作。看起来，开发人员不再关心全局存储库了。但是全局内存仍然很慢，所以正确的访问模式很重要。现在的重点是尽可能多地使用/重用 L2。我的问题是，如何？如果我需要一些详细的信息，L2 是如何工作的，以及我应该如何组织和访问全局内存，例如每个线程 100-200 个元素数组，我将不胜感激。

Answer 1

L2 缓存在某些方面有所帮助，但它并不能消除对全局内存的合并访问的需要。简而言之，合并访问意味着对于给定的读取(或写入)指令，warp 中的各个线程正在读取(或写入)全局内存中相邻的连续位置，最好在 128 字节边界上对齐为一个组.这将导致对可用内存带宽的最有效利用。

在实践中，这通常不难实现。例如:

int idx=threadIdx.x + (blockDim.x * blockIdx.x);
int mylocal = global_array[idx];

假设 global_array，将提供跨经线中所有线程的合并(读取)访问权限在全局内存中使用 cudaMalloc 以普通方式分配。这种类型的访问使可用内存带宽的使用率达到 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对这个话题很有启发性。您可能还想阅读 Device Memory Access section的 CUDA C Programming Guide .