CUDA编程 基础与实践 学习笔记（六）

6.1 CUDA的内存组织
内存存在一种层级结构，大容量的一般latency高，小容量的一般latency低。这一点在计算机组成原理中学到过。
这里有个疑问，表中的局部内存是放在芯片外部的，这里怎么画在了内部？还是我没有理解清楚？
答：局部变量属于全局变量的一部分所以在芯片外。当寄存器放不下了，就会放在局部变量中，由寄存器做判断。

6.2 集中不同类型的内存：
（1）全局内存，核函数所有线程都能访问，就是显卡的显存，低速高延迟(速度慢是相对于数据处理而言的)。负责host->device, device->device的数据传递，用cudaMalloc，cudaMemcpy进行操作，可读可写。生命周期由主机端决定(之前分配内存时，使用的是**指针，存储在host上)，用cudaFree()进行释放。之前讨论的一直是动态分配内存，也可以静态分配内存，所占用内存大小在动态分配时就确定，必须在主机和设备函数外部定义。__device__ T x[N]。在核函数中可以直接访问，在主机函数中要做移动，用cudaMemcpyFromSymbol, cudaMemcpyToSymbol()(很自然，因为这是定义在device上的)

(2) 常量内存， 仅64KB，有缓存。可见范围生命周期同全局内存，read only，访存速度快。在核函数外使用__constant__定义变量，用cudaMemcpyToSymbol()复制，给核函数传递参数最多只能在核函数内使用4KB(注意理解常量内存有缓存的含义)常量内存。

(3) 纹理内存和表面内存。名字蛮有趣。

（4）寄存器。速度最快，容量最小，可读可写，仅能被一个线程可见，生命周期与线程一致。一般不加限定的变量、数组有可能放在寄存器中（有些放在局部内存中）。gridDim, blockDim这些内建变量都在特殊寄存器中，访问很高效。
会被放在寄存器中。

（5）局部内存。属于全局变量（这就解释了为什么在片外。）刚刚提到，核函数中不加限定符的变量可能放在寄存器中，也可能放在局部变量中。寄存器中放不下的变量，索引值不能在编译时就确定的数组，都有可能放在局部内存中，由编译器来做判断。延迟很高，每个线程最多使用512KB的局部内存。

（6）共享内存。在芯片上，读写速度仅次于寄存器。但与寄存器不同的是，这玩意在一定程度上共享（相当于进程的内存？进程内所有线程都可以享有），每一个线程都有一个共享内存副本，不同线程的共享内存的值可以不同，但不能跨线程访问。主要作用是减少对全局内存的访问。

（7）缓存。用来缓存全局变量和局部变量的访问，减少延迟。不可编程。

6.3 SM 及其占有率

https://blog.csdn.net/qq_41554005/article/details/119757653感觉比书上介绍的清楚。。
这个Blog回答了我的一个疑问：GPU的计算能力指的是SM的功能，而非数量。架构越新，SM的功能越强。理论上GPU的计算能力决定驻留在SM上的线程块的最大数量。（可以理解为SM定义最多能运行的线程块数量）。
SM中包含流处理器，streaming processor,也就是cuda core. 一个SM中包含几十上百个SP，每个只运行一个线程，有自己的寄存器集，相当于CPU中的执行单元。GPU之所以可以并行计算，是因为GPU的体系结构有很多个SP同时进行处理，通过单指令多线程来实现(原来概念辨析不清楚，以为可以在全局变量的尺度下(例如16GB)进行矩阵运算，感觉还蛮壮观的。原来是靠几百个SP来并行计算的，嗯…也行吧，不如我设想的壮观，hhh)。虽然SP这个术语有点弱化了，直接用thread来代替，但对于初学者来说，SP是一个硬件上的概念，有助于软硬件的区分与对应。CUDA core是单进度的，有双精度的dp unit, double precision unit.

既然并行计算是靠单指令多线程的SM来实现的，那么SM中的SP占有率就及其重要。一般不小于某个值， 才能获得比较高的性能。

1. block数量要取32的整倍数。因为线程束的大小为32。不完整的线程束与完整的花费的时间一样。在寄存器和共享内存使用量很少的情况下，只要$N_t/N_b$能够整除，就能得到100%的利用率。

疑问：怎么看线程占了多少个寄存器呢？

我的2060：

这一章偏科普，学到了很多概念，对GPU内部结构有了新的认识，很有趣~