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

Question

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

4

从CUDA计算能力2.0（Fermi）开始，全局内存访问通过768 KB L2缓存进行。开发者似乎不再关心全局内存块了。但是全局内存仍然非常慢，因此正确的访问模式很重要。现在的问题是尽可能多地使用/重复使用L2。我的问题是，如何做到这一点？我将感激一些详细信息，包括L2的工作原理以及如果我需要每个线程处理100-200个元素数组时应该如何组织和访问全局内存。

- Andrew

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

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网络研讨会来获得更深入的指导。例如，全局内存使用和策略网络研讨会（以及幻灯片）或CUDA共享内存和缓存网络研讨会对于这个主题是有益的。您还可以阅读CUDA C编程指南的设备内存访问部分。