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只访问不多于一个内存地址，那么只需要一次内存传输就可以了，否则需要多次传输，因此会降低内存带宽的使用。