4

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

4

1 回答 1

9

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。

Some advice I would offer is to not assume that the L2 cache just fixes sloppy memory accesses. The GPU caches are much smaller than equivalent caches on CPUs, so it's easier to get into trouble there. A general piece of advice is simply to code as if the caches were not there. Rather than CPU oriented strategies like cache-blocking, it's usually better to focus your coding effort on generating coalesced accesses and then possibly make use of shared memory in some specific cases. Then for the inevitable cases where we can't make perfect memory accesses in all situations, we let the caches provide their benefit.

You can get more in-depth guidance by looking at some of the available NVIDIA webinars. For example, the Global Memory Usage & Strategy webinar (and slides ) or the CUDA Shared Memory & Cache webinar would be instructive for this topic. You may also want to read the Device Memory Access section of the CUDA C Programming Guide.

于 2012-12-12T08:07:48.050 回答