CUDA Compute Capability 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];

将在warp中的所有线程上提供合并(读取)访问,假设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网络研讨会,您可以获得更深入的指导.例如,全球内存使用情况及策略的网络研讨会(和幻灯片 )或CUDA共享存储器和缓存网络研讨会将能供此主题.您可能还想阅读" CUDA C编程指南"的" 设备内存访问"部分.