CUDA on Platform 学习笔记6--多种CUDA存储单元

本文是参加2022 CUDA on Platform 线上训练营学习笔记，感谢NVIDIA各位老师的精彩讲解!

前面介绍简单介绍过多种memory的联系和区别，本节进行详解，下图是所有memory和thread，block，grid的关系图，前面第4节已介绍过，不再赘述。

1. Registers:
   寄存器是GPU最快的memory，没有之一， kernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中。
   1）寄存器变量是每个线程私有的，一旦thread执行结束，寄存器变量就会失效。
   2） 寄存器是稀有资源。(省着点用，能让更多的block驻留在SM中，运行的block数量越多，并行化越好，增加Occupancy，执行效率越高)。
   3） --maxrregcount 可以设置大小
   4） 不同设备架构，数量不同
2. Shared Memory:
   用__shared__修饰符修饰的变量存放在shared memory:
   1）On-chip
   2） 拥有高的多bandwidth和低很多的latency，share memory 速度接近register
   3）同一个Block中的线程共享一块Shared Memory
   4） __syncthreads()同步，share memory是多个线程共享，因此需要同步
   5）比较小，要节省着使用，不然会限制活动warp的数量
3. Local Memory:
   有时候，Registers不够了，就会用Local Memory来替代。但是，更多在以下情况，会使用Local
   Memory:
   1）无法确定其索引是否为常量的数组。
   2）会消耗太多寄存器空间的大型结构或数组。
   3） 如果内核使用了多于可用寄存器的任何变量(这也称为寄存器溢出)
   4） --ptxas -options=-v
4. Constant Memory:
   固定内存空间驻留在设备内存中，并缓存在固定缓存中(constant cache):
   1）constant的范围是全局的，针对所有kernel
   2）在同一个编译单元，constant对所有kernel可见。
   3）kerne[只能从constant Memory读取数据，因此其初始化必须在host端使用下面的function调用:cudaError_t cudaMemcpyToSymbo[【const voidsymbo[，const voidsrc,size_t count);
   4）当一个warp中所有thread都从同一个Memory地址读取数据时，constant Memory表现会非常好，会触发广播机制。
5. Texture Memory
   1）Texture Memory驻留在device memory 中，并且使用一个只读cache。它实际上也是global memory的一块，但它有自己专有的只读cache。
   2）纹理内存也是缓存在片上的，因此一些情况下相比从芯片外的DRAM上获取数据，纹理内存可以通过减少内存请求来提高带宽。
   3）它是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。在某个计算应用程序中，这意味着一个线程读取的位置可能与邻近线程的读取位置“非常接近”。
   
   从数学的角度，上图中的4个地址并非连续的，在一般的CPU缓存中，这些地址将不会缓存。但由于GPU纹理缓存是专门为了加速这种访问模式而设计的，因此如果在这种情况中使用纹理内存而不是全局内存，那么将会获得性能的提升。
6. Global Memory
   空间最大，latency最高，GPU最基础的memory
   1）驻留在device memory 中
   2）memory transaction对齐，合并访存。
   由于Global Memory合并访存的特性，每次会读取32个数据，若是只使用一个数据，其他的数据其实是浪费掉了。如果后面的数据被相邻的线程使用，效率就会很高，因此，连续的线程读取连续的数据（一行）效率最高。
   注意和cpu指针读取二维数组区分，CPU Cache从内存中抓取一般都是整个数据块，所以它的物理内存是连续的，几乎都是同行不同列，而如果内循环以列的方式进行遍历的话，将会使整个缓存块无法被利用，而不得不从内存中读取数据，而从内存读取速度是远远小于从缓存中读取数据，因此此时按行遍历效率高。而gpu是多个线程并行，因此对于连续的线程按行就是读取连续数据，对于每一个线程是按列读取。具体可以看下图示例。
   左图是每个线程按行读取连续数据，连续的线程T0–T4按列读取数据（不连续），右图是每个线程按列读取数据，连续的线程T0–T4按行读取连续数据，右图的效率更高。