cuda的shared momery

CUDA SHARED MEMORY

在global Memory部分,数据对齐和连续是很重要的话题,当使用L1的时候,对齐问题可以忽略,但是非连续的获取内存依然会降低性能。依赖于算法本质,某些情况下,非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。

GPU上的memory有两种:

· On-board memory

· On-chip memory

global memory就是一块很大的on-board memory,并且有很高的latency。而shared memory正好相反,是一块很小,低延迟的on-chip memory,比global memory拥有高得多的带宽。我们可以把他当做可编程的cache,其主要作用有:

· An intra-block thread communication channel 线程间交流通道

· A program-managed cache for global memory data可编程cache

· Scratch pad memory for transforming data to improve global memory access patterns

本文主要涉及两个例子作解释:reduction kernel,matrix transpose kernel。

shared memory(SMEM)是GPU的重要组成之一。物理上,每个SM包含一个当前正在执行的block中所有thread共享的低延迟的内存池。SMEM使得同一个block中的thread能够相互合作,重用on-chip数据,并且能够显著减少kernel需要的global memory带宽。由于APP可以直接显式的操作SMEM的内容,所以又被称为可编程缓存。

由于shared memory和L1要比L2和global memory更接近SM,shared memory的延迟比global memory低20到30倍,带宽大约高10倍。

 

当一个block开始执行时,GPU会分配其一定数量的shared memory,这个shared memory的地址空间会由block中的所有thread 共享。shared memory是划分给SM中驻留的所有block的,也是GPU的稀缺资源。所以,使用越多的shared memory,能够并行的active就越少。

关于Program-Managed Cache:在C语言编程里,循环(loop transformation)一般都使用cache来优化。在循环遍历的时候使用重新排列的迭代顺序可以很好利用cache局部性。在算法层面上,我们需要手动调节循环来达到令人满意的空间局部性,同时还要考虑cache size。cache对于程序员来说是透明的,编译器会处理所有的数据移动,我们没有能力控制cache的行为。shared memory则是一个可编程可操作的cache,程序员可以完全控制其行为。

Shared Memory Allocation

我们可以动态或者静态的分配shared Memory,其声明即可以在kernel内部也可以作为全局变量。

其标识符为:__shared__。

下面这句话静态的声明了一个2D的浮点型数组:

__shared__ float tile[size_y][size_x];

如果在kernel中声明的话,其作用域就是kernel内,否则是对所有kernel有效。如果shared Memory的大小在编译器未知的话,可以使用extern关键字修饰,例如下面声明一个未知大小的1D数组:

extern __shared__ int tile[];

由于其大小在编译器未知,我们需要在每个kernel调用时,动态的分配其shared memory,也就是最开始提及的第三个参数:

kernel<<<grid, block, isize * sizeof(int)>>>(…)

应该注意到,只有1D数组才能这样动态使用。

Memory Banks

为了获得高带宽,shared Memory被分成32(对应warp中的thread)个相等大小的内存块,他们可以被同时访问。不同的CC版本,shared memory以不同的模式映射到不同的块(稍后详解)。如果warp访问shared Memory,对于每个bank只访问不多于一个内存地址,那么只需要一次内存传输就可以了,否则需要多次传输,因此会降低内存带宽的使用。

 

Published by

风君子

独自遨游何稽首 揭天掀地慰生平

发表回复

您的电子邮箱地址不会被公开。 必填项已用 * 标注