2012-12-12 40 views
4

从CUDA Compute Capability 2.0(Fermi)开始,全局内存访问通过768 KB二级高速缓存运行。看起来,开发商不再关心全球内存银行。但全球记忆仍然非常缓慢,所以正确的访问模式非常重要。现在重点是尽可能地使用/重用L2。我的问题是,怎么样?我会感谢一些详细的信息,L2如何工作以及如何组织和访问全局内存(如果需要),例如,每个线程100-200个元素数组。CUDA计算能力2.0。全局内存访问模式

回答

9

二级高速缓存在某些方面有所帮助,但并不排除对全局内存进行合并访问的需要。简而言之,合并访问意味着对于给定的读取(或写入)指令,变形中的单个线程正在读取(或写入)全局内存中邻近的连续位置,最好是在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缓存行为类似。缓存行的概念与缓存服务请求的行为类似。

费米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 GuideDevice Memory Access section

+0

除了这个优秀的答案,费米后硬件还有一个额外的专用只读L1(与正常读写L1的大小相同)。这意味着如果有任何理由要关注这些“小细节”,现在就是了。 – Damon

+0

@RobertCrovella很好的答案,但我有一个关于聚结阅读的问题。对于全局内存的聚结读取,我们必须使用__syncthreads();'。在上面的例子中'int local1 = global_array [idx];'后面会跟着'__syncthreads();'。问题是,如果我有多个数组像'int local1 = global_array1 [idx]; int local2 = global_array2 [idx]; int local3 = global_array3 [idx];'在所有这些定义之后,我会用一个'__syncthreads();来读取聚结的吗?谢谢。 – BugShotGG

+0

合并读取和'__syncthreads()'没有任何关系。 –