CUDA_C_NOTES [5]

第5章 共享内存和常量内存

• 了解数据在共享内存中是如何被安排的
• 掌握从二维共享内存到线性全局内存的索引转换
• 解决不同访问模式中存储体中的冲突
• 在共享内存中缓存数据以减少对全局内存的访问
• 使用共享内存避免非合并全局内存的访问
• 理解常量缓存和只读缓存之间的差异
• 使用线程束洗牌指令编程

5.1 CUDA共享内存概述

GPU中有两种类型的内存:

• 板载内存: 全局内存是较大的板载内存，具有相对较高的延迟。
• 片上内存: 共享内存是较小的片上内存，具有相对较低的延迟，并且共享内存可以提供比全局内存高得多的带宽

共享内存通常的用途有:

• 块内线程通信的通道
• 用于全局内存数据的可编程管理的缓存
• 高速暂存存储器，用于转换数据以优化全局内存访问模式

5.1.1 共享内存

共享内存（shared memory，SMEM）是GPU的一个关键部件。物理上，每个SM都有一个小的低延迟内存池，这个内存池被当前正在该SM上执行的线程块中的所有线程所共享。(共享内存就是SM上的一块低延迟内存池)

共享内存使同一个线程块中的线程能够互相协作，便于重用片上数据，并可以大大降低核函数所需的全局内存带宽。由于共享内存中的内容是由应用程序显式管理的，所以它通常被描述为可编程管理的缓存。

当每个线程块开始执行时，会分配给它一定数量的共享内存。这个共享内存的地址空间被线程块中所有的线程共享。它的内容和创建时所在的线程块具有相同生命周期。每个线程束发出共享内存访问请求。在理想的情况下，每个被线程束共享内存访问的请求在一个事务中完成。最坏的情况下，每个共享内存的请求在32个不同的事务中顺序执行。如果多个线程访问共享内存中的同一个字，一个线程读取该字后，通过多播把它发送给其他线程。

共享内存被SM中的所有常驻线程块划分，因此，共享内存是限制设备并行性的关键资源。一个核函数使用的共享内存越多，处于并发活跃状态的线程块就越少。

可编程管理的缓存

共享内存是一个可编程管理的缓存。当数据移动到共享内存中以及数据被释放时,我们对它有充分的控制权。由于在CUDA中允许手动管理共享内存,所以通过在数据布局上提供更多的细粒度控制和改善片上数据的移动,使得对应用程序代码进行优化变得更简单了

5.1.2 共享内存分配

有多种方法可以用来分配或声明由应用程序请求所决定的共享内存变量。可以静态或动态地分配共享内存变量。在CUDA的源代码文件中,共享内存可以被声明为一个本地的CUDA核函数或是一个全局的CUDA核函数。CUDA支持一维、二维和三维共享内存数组的声明。

共享内存变量用下列修饰符进行声明: __shared__

如果在核函数中进行声明,那么这个变量的作用域就局限在该内核中。如果在文件的任何核函数外进行声明,那么这个变量的作用域对所有核函数来说都是全局的。