【CUDA】Shared Memory

共享内存(SMEM)是 GPU 的一个关键部件。物理上，每个 SM 都有一个小的低延迟内存池，这个内存池被当前正在该 SM 上执行的线程块中的所有线程所共享。共享内存使同一个线程块中的线程能够互相协作，便于重用片上数据，并可以大大降低核函数所需的全局内存带宽。由于共享内存中的内容是由应用程序显式管理的，所以其通常被描述为可编程管理的缓存

内存层次结构如下图所示，全局内存的所有加载和存储请求都要经过二级缓存，这是 SM 单元之间数据统一的基本点。相较于二级缓存和全局内存，共享内存和一级缓存在物理上更接近 SM。因此，共享内存相对于全局内存而言，延迟要低大约 20~30 倍，而带宽高其大约 10 倍

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

为什么说共享内存是一个可编程管理的缓存？

在C语言中，循环转换是一种常见的缓存优化方法。通过重新安排迭代顺序，循环转换可以在循环遍历的过程中提高缓存局部性。在算法层面上，在考虑缓存大小的同时，可以手动调整循环，以实现更好的空间局部性

缓存对程序而言是透明的，编译器可以处理所有的数据移动，而程序员不能控制缓存的释放。但是当数据移动到共享内存中以及数据被释放时，程序员对其有充分的控制权。由于在CUDA中允许手动管理共享内存，所以通过在数据布局上提供更多的细粒度控制和改善片上数据的移动，使得对应用程序代码进行优化变得更简单了

二、共享内存分配

可以动态或者静态的分配 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数组才能这样动态使用

三、存储体和访问模式

3.1 内存存储体(Memory Banks)

为了获得高内存带宽，共享内存被分为 32 个（对应warp中的thread）同样大小的内存模型，被称为存储体(bank)，可以被同时访问。共享内存是一个一维地址空间。根据GPU的计算能力，共享内存的地址在不同模式下会映射到不同的存储体中。若通过 warp 发布共享内存加载或存储操作，且在每个存储体上只访问不多于一个的内存地址，那么只需要一次内存传输就可以了，否则需要多次传输，因此会降低内存带宽的使用

3.2 存储体冲突(Bank Conflict)

当多个地址请求落在同一个 bank 中就会发生 bank conflict，从而导致请求多次执行。硬件会把这类请求分散到尽可能多的没有 conflict 的那些传输操作中，降低有效带宽的因素是被分散到的传输操作个数

warp有三种典型的获取 shared memory 的模式：

并行访问（Parallel access）：多个地址访问多个存储体
串行访问（Serial access）：多个地址访问同一个存储体
广播访问（Broadcast access）：单一地址读取单一存储体

Parallel access 是最通常的模式，这个模式一般暗示，一些（也可能是全部）地址请求能够被一次传输解决。最佳情况是，当每个地址都位于一个单独的存储体时，执行无冲突的共享内存访问

Serial access是最坏的模式，当多个地址属于同一个存储体时，必须以串行的方式进行请求，若 warp 中的 32 个 thread 都访问了同一个 bank 中的不同位置，那就是 32 次单独的请求，而不是同时访问了

Broadcast access，warp 中所有的线程都读取同一存储体中相同的地址，然后传输结果会广播给所有发出请求的thread，这样的话就会导致带宽利用率低

最优情况访问图示

下图是随机访问，同样没有conflict：

某些thread访问到同一个bank的情况

若线程访问 bank 内的相同地址，则无冲突广播访问
若线程访问 bank 内的不同地址，则进行 bank 冲突访问

3.3 访问模式

共享内存存储体的宽度规定了共享内存地址于共享内存存储体的对应关系。内存存储体的宽度随设备计算能力的不同而变化。有两种不同的存储体宽度：

计算能力2.x的设备中为4字节（32位）
计算能力3.x的设备中为8字节（64位）

存储体宽度为4字节

对于Fermi设备，存储体的宽度是 32 位并且有 32 个存储体。每个存储体在每两个时钟周期内都有32位的带宽。连续的 32 位字映射到连续的存储体中，因此，从共享内存地址到存储体索引的映射可以按下面公式进行计算：

存储体索引 = （字节地址 ÷ 4 bytes/bank） % 32 banks

字节地址除以4转换为一个4字节字索引，然后进行模32操作，将4字节字索引转换为存储体索引

下图所示的上部显示了在 Fermi 设备中从字节地址到字索引的映射，下部显示了从字索引到存储体索引的映射

注意：存储体成员线束相差 32 * 4 个字节。邻近的字被分到不同的存储体中，以最大限度地提高线程束中可能的并发访问数量

当来自相同线程束中的两个线程访问相同的地址时，大概率不会发生存储体冲突。在这种情况下，对于读访问，这个字被广播到请求的线程中；对于写访问，这个 4字节只由其中一个线程写入

存储体宽度为8字节

对于Kepler设备，共享内存有32个存储体，有以下两种地址模式：64位模式和32位模式。在64位模式下，连续的64位字映射到连续的存储体中，在每时钟周期内每个存储体都有64位的带宽，从共享内存地址到存储体索引的映射可以按以下公式来进行计算：

存储体索引 = (字节地址 ÷ 8 bytes/bank) % 32 banks

两个 thread 访问同一个 64-bit 中的任意 word 也不会导致 bank conflict，因为一次 64-bit（bank带宽64bit/cycle）的读就可以满足请求了。在相同的访问模式下，64-bit 模式一般比 32-bit 模式更少碰到 bank conflict

在 32 位模式下，连续的 32 位字映射到连续的存储体中。然而，因为 Kepler 在每个时钟周期内都有 64 位带宽，在同一存储体中访问两个 32 位字并不总意味着重操作。在单一的时钟周期内读 64 位并只将 32 位请求传输给每个线程，这是有可能的

下图显示了在 32 位模式下从字节地址到存储体索引上的映射。上部的图是字节地址和 4 字节索引标记的共享内存。下部的图显示了从4字节索引到存储体索引的映射。虽然 word 0 和 word 32 都在 bank 0 中，但是在相同的内存请求中读取这两个字不会产生存储体冲突

下图显示了在 64 位模式下无冲突访问的一种情况，在这种情况下，每个线程访问不同的存储体

下图显示了在64位模式下无冲突访问的另一种情况，在这种情况下，两个线程访问相同存储体中的字和相同的8字节字

下图展示了一个双向存储体冲突，在这种情况下，两个线程访问同一个存储体，但地址落在两个不同的8字节字中

3.4 内存填充

内存填充是避免存储体冲突的一种方法，下图通过一个简单的例子来说明内存填充。假设只有5个共享内存存储体。若所有线程访问 bank 0 的不同地址，那么会发生一个五向的存储体冲突。解决这种存储体冲突的一个方法是在每 N 个元素之后添加一个元素，N 是存储体的数量。这就改变了从字到存储体的映射。如下图右侧所示，由于填充，之前所有属于 bank 0 的元素，现在被传播到了不同的存储体中

填充的内存不能用于数据存储，其唯一的作用是移动数据元素，以便将原来属于同一个存储体中的数据分散到不同的存储体中。这样，线程块可用的总的共享内存的数量将减少。填充之后，还需要重新计算数组索引以确保能访问到正确的数据元素。虽然 Fermi 和 Kepler 都有32个存储体，但存储体宽度不同，在这些不同的架构上填充共享内存时，必须要小心。Fermi架构中的某些内存填充模式可能会导致Kepler中的存储体冲突

3.5 访问模式配置

Kepler 设备支持 4字节到 8字节的共享内存访问模式。默认是 4字节模式。可采用以下的CUDA运行时API函数查询访问模式：

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);

结果返回到pConfig中。返回的存储体配置可以是下列值中的一个：

cudaSharedMemBankSizeFourByte
cudaSharedMemBankSizeEightByte

在可配置共享内存存储体的设备上，可以使用以下功能设置一个新的 bank 大小：

cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);

支持的存储体配置为：

cudaSharedMemBankSizeDefault
cudaSharedMemBankSizeFourByte
cudaSharedMemBankSizeEightByte

在不同的核函数启动之间更改共享内存配置可能需要一个隐式的设备同步点。更改共享内存存储体的大小不会增加共享内存的使用量，也不会影响核函数的占用率，但对性能可能有重大影响。一个大的存储体可能为共享内存访问产生更高的带宽，但是可能会导致更多的存储体冲突，这取决于应用程序中共享内存的访问模式

四、配置共享内存量

每个 SM 都有 64KB 的片上内存，共享内存和一级缓存共享该硬件资源。CUDA为配置一级缓存和共享内存的大小提供了两种方法：

按设备进行配置
按核函数进行配置

按设备进行配置

使用下述的运行时函数，可以为在设备上启动的核函数配置一级缓存和共享内存的大小：

cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);

参数cacheConfig指明，在当前的CUDA设备上，片上内存是如何在一级缓存和共享内存间进行划分，所支持的缓存配置参数如下所示：

cudaFuncCachePreferNone // no preference(default)
cudaFuncCachePreferShared // prefer 48KB shared memory and 16KB L1 cache
cudaFuncCachePreferL1 // prefer 48KB L1 cache and 16KB shared memory
cudaFuncCachePreerEqual // perfer 32KB L1 cache and 32KB shared memory

哪种模式更好，这取决于在核函数中使用了多少共享内存。典型情况如下：当核函数使用较多的共享内存时，倾向于更多的共享内存；当核函数使用更多的寄存器时，倾向于更多的一级缓存。若核函数使用了大量的共享内存，那么配置 48KB 的共享内存能实现较高的占用率核更好的性能；若核函数使用了少量的共享内存，那么应该为一级缓存配置 cacheConfig 参数为48KB

按核函数进行配置

CUDA运行时会尽可能使用请求设备的片上内存配置，但是若需要执行一个核函数，可自由选择不同的配置，每个核函数的配置可以覆盖设备范围的设置

使用以下运行时函数进行设置：

cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCacheca cheConfig);

核函数使用情况是由核函数指针func指定的，启动一个不同优先级的内核比启动有最近优先级设备的内核更可能会导致隐式设备同步。对于每个核，只需调用一次这个函数。每个核函数启动时，片上内存中的配置不需要重新设定

L1 cache 与 shared memory

虽然一级缓存和共享内存位于相同的片上硬件，但在某些方面它们却不太相同，共享内存是通过32个存储体进行访问的，而一级缓存则是通过缓存行进行访问的。使用共享内存，对存储内容和存放位置有完全的控制权，而使用一级缓存，数据删除工作是由硬件完成

一般情况下，GPU 缓存的行为比 CPU 缓存的行为更难以理解。GPU 使用不同的启发式算法删除数据。在 GPU 上，数百个线程共享相同的一级缓存，数千个线程共享相同的二级缓存。因此，数据删除在 GPU 上可能会发生得更频繁而且更不可预知，使用 GPU 共享内存不仅可以显示管理数据而且还可以保证 SM 的局部性

五、同步

并行线程间的同步是所有并行计算语言的重要机制。正如它名字所暗示的，共享内存可以同时被线程块中的多个线程访问。当不同步的多个线程修改同一个共享内存地址时，将导致线程内的冲突。CUDA提供了几个运行时函数来执行块内同步。同步的两个基本方法如下所示：障碍和内存栅栏。在障碍中，所有调用的线程等待其余调用的线程到达障碍点。在内存栅栏中，所有调用的线程必须等到全部内存修改对其余调用线程可见时才能继续执行。然而，在学习CUDA的块内障碍点和内存栅栏之前，理解CUDA调用的弱排序顺序内存模型是非常重要的

5.1 弱排序内存模型

现代内存架构有一个宽松的内存模型。内存访问不一定按照在程序中出现的顺序进行执行。CUDA采用弱排序内存模型从而优化了更多激进的编译器

GPU线程在不同内存（如共享内存、全局内存、锁页主机内存或对等设备的内存）中写入数据的顺序，不一定和这些数据在源代码中访问的顺序相同。一个线程的写入顺序对其他线程可见时，其可能和写顺序被执行的实际顺序不一致。若指令之间是相互独立的，线程从不同内存中读取数据的顺序和读指令在程序中出现的顺序不一定相同。为了显式地强制程序以一个确切地顺序执行，必须在应用程序中插入内存栅栏和障碍。这是保证与其他线程共享资源的核函数行为正确的唯一途径

5.2 障碍

在CUDA中，障碍只能在同一个线程块的线程间执行。在核函数中，可以通过调用下面地函数来指定一个障碍点：

void __syncthreads();

__synthreads作为一个障碍点来发挥作用，要求块中的线程必须等待直到所有线程都到达该点。__synthreads还确保在障碍点之前，被这些线程访问的所有全局和共享内存对同一块中的所有线程都可见。__synthreads用于协调同一块中线程间的通信

当块中的某些线程访问共享内存或全局内存中的同一地址时，会有潜在问题(写后读、读后写、写后写)，这将导致在那些内存位置产生未定义的应用程序行为和未定义的状态，可以通过利用冲突访问间的同步线程来避免这种情况

在条件代码中使用__synthreads时，必须要小心，若一个条件能保证对整个线程块进行同等评估，则它是调用__synthreads的唯一有效条件。否则执行很可能会挂起或产生意想不到的的问题。如，下面的代码可能会导致块中的线程无限期的等待对方，因为块中的所有线程没有达到相同的障碍点

if(threadID % 2 == 0)
{
   __synthreads();
}
else
{
   __synthreads();
}

不同线程块之间的同步并不直接支持，线程块可能会以任何顺序、并行、串行的顺序在 SM 上执行。若一个CUDA核函数要求跨线程块全局同步，那么通过在同步点分割核函数并执行多个内核启动可能会达到预期的效果。因为每个连续的内核启动必须等待之前的内核启动完成，所以这会产生一个隐式全局障碍

5.3 内存栅栏

内存栅栏的功能可确保栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。根据所需范围，有3种内存栅栏：块、网格或系统。通过以下固有函数可以在线程块内创建内存栅栏：

void __threadfence_block();

_threadfence_block保证了栅栏前被调用线程产生的对共享内存和全局内存的所有写操作对栅栏后的同一块中的其他线程都是可见的。内存栅栏不执行任何线程同步，所以对于一个块中的所有线程来说，没有必要执行这个指令

使用下面的固有函数来创建网格级内存栅栏：

void __threadfence();

__threadfence 挂起调用的线程，直到全局内存中所有的写操作对相同网格内的所有线程都可见

使用下面的函数可以跨系统(包括主机和设备)设置内存栅栏：

void __threadfence_system();

__threadfence_system 挂起调用的线程，以确保该线程对全局内存、锁页主机内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程是可见的

六、Volatile修饰符

在全局或共享内存中使用 volatile修饰符声明一个变量，可以防止编译器优化，编译器优化可能会将数据暂时缓存在寄存器或本地内存中。当使用 volatile修饰符时，编译器假定任何其他线程在任何时间都可以更改或使用该变量的值，强制每次都到 global 或者 shared Memory 中去读取其绝对有效值

七、全局内存与共享内存

GPU全局内存常驻在设备内存(DRAM)上，比GPU的共享内存访问慢得多。相较于DRAM，共享内存有以下几个特点：DRAM比其高20~30倍的延迟；比DRAM大10倍的带宽。共享内存的访问粒度也比较小。而DRAM的访问粒度可以是32个字节或128个字节，共享内存的访问粒度如下：Fermi架构：4字节存储体宽；Kepler架构：8字节存储体宽