3、CUDA编程模型（上）：CUDA架构概览、线程层次结构（Grid/Block/Warp）、内存层次结构（Global/Shared/Local/Register）
好，咱们今天正式进入CUDA编程模型。说实话，我当年刚接触CUDA时，最头疼的就是搞不清Grid、Block、Warp这些概念到底怎么对应到硬件上。你想想看，一个程序要跑在几千个核上，怎么分工？怎么协作？数据放哪里？这些问题不搞清楚，写出来的代码要么跑不起来，要么跑得比CPU还慢。
这一节，咱们先把CUDA的架构骨架搭起来。我会结合我这些年踩过的坑，帮你把线程层次和内存层次这两个核心模型彻底吃透。
3.1 CUDA架构概览：从GPU到CUDA的抽象
先看一张图，这是我个人习惯用来理解CUDA整体架构的方式。它把硬件和软件模型对应了起来。
这张图我建议你多看几遍。左侧是咱们写代码时看到的抽象层次，右侧是GPU硬件实际干活的样子。说白了，CUDA编程模型就是给咱们一个「骗自己」的接口——让你觉得好像有无数个线程在并行跑，而硬件层面其实是在用SIMT（单指令多线程）的方式批量执行。
核心要点：CUDA的编程模型是「单程序多数据」（SPMD）的变体。你写一个kernel函数，然后让成千上万个线程去执行它，每个线程处理不同的数据。
3.2 线程层次结构：Grid / Block / Warp
好，咱们一层层拆开来看。
3.2.1 Grid（网格）
Grid是最大的线程组织单位。当你调用一个kernel时，实际上就是启动了一个Grid。一个Grid包含若干个Block。
// 定义一个Grid，包含2x3=6个Block
dim3 gridDim(2, 3);
dim3 blockDim(4, 4);
// 启动kernel
myKernel<<<gridDim, blockDim>>>(d_data);这里有个细节：Grid的维度可以是1D、2D或3D。我个人习惯用1D处理数组，2D处理图像，3D处理体数据。你想想看，处理一张1920x1080的图片，用2D Grid天然就对应了像素坐标，多直观。
3.2.2 Block（线程块）
Block是Grid里的子单位。同一个Block内的线程可以协作——它们能共享数据，也能同步。不同Block之间的线程是独立的，不能直接通信。
这里有个关键限制：一个Block最多1024个线程（取决于GPU架构）。为什么？因为Block内的线程要共享SM上的资源，比如寄存器和共享内存。资源就那么多，线程太多就分不过来了。
我曾经踩过的坑：刚开始写CUDA时，我习惯把Block设成(32, 32)，也就是1024个线程。结果发现某些老显卡（Compute Capability 2.x）只支持512个线程/Block。程序直接崩了，查了半天才找到原因。所以，写代码前先查一下目标GPU的规格。
3.2.3 Warp（线程束）
Warp是GPU执行的最小单位。一个Warp包含32个线程。这32个线程在硬件上执行相同的指令，但处理不同的数据——这就是SIMT的精髓。
为什么是32？这是NVIDIA的设计选择。32个线程一起取指、译码、执行，能最大化利用指令带宽。你想想看，如果一次只取一条指令给一个线程，那指令缓存得被频繁访问，效率就低了。
| 概念 | 软件/硬件 | 典型大小 | 关键特性 |
|---|---|---|---|
| Grid | 软件 | 由Block数量决定 | 一次Kernel调用对应一个Grid |
| Block | 软件 | 最多1024线程 | 同一Block内可共享内存、可同步 |
| Warp | 硬件 | 32线程 | SIMT执行单元，Warp内线程同步执行 |
| Thread | 软件 | 1 | 最小的执行上下文 |
个人经验：Block的维度最好设成32的倍数。为什么？因为Warp是32个线程一组。如果你的Block有64个线程，那正好2个Warp。如果有96个，那就是3个Warp。但如果Block有100个线程，那最后一个Warp只有4个线程在干活，剩下的28个在「空转」，浪费了计算资源。
3.3 内存层次结构：Global / Shared / Local / Register
搞清楚了线程怎么组织，接下来就是数据放哪的问题。GPU的内存层次比CPU更复杂，但也更讲究。用对了，性能起飞；用错了，性能惨不忍睹。
3.3.1 寄存器（Register）
寄存器是离计算单元最近的内存。每个线程都有自己的寄存器，访问延迟几乎为0。但寄存器数量有限——每个SM的寄存器总数是固定的（比如65536个），分给所有线程。
__global__ void kernel() {
int tid = threadIdx.x; // tid存在寄存器里
float a = 3.14f; // a也在寄存器里
float b = a * tid; // 计算也在寄存器里完成
}这里有个坑：如果你在kernel里声明了太多局部变量，编译器会「溢出」到Local Memory。Local Memory虽然名字带「Local」，实际上存在显存里，速度慢得多。
我曾经遇到过：一个kernel里用了20多个float变量，结果性能比预期慢了10倍。一查，寄存器不够用，变量被溢出到Local Memory了。解决办法是减少每个线程的变量数，或者减少Block大小（让更多寄存器分给每个线程）。
3.3.2 局部内存（Local Memory）
Local Memory是寄存器的「备胎」。当寄存器不够用时，编译器自动把变量放到Local Memory。它在物理上位于显存（Global Memory）中，但有L1/L2缓存加速。
访问Local Memory的延迟大约是寄存器的100倍。所以，尽量避免寄存器溢出。
3.3.3 共享内存（Shared Memory）
共享内存是CUDA编程的「杀手锏」。它位于SM内部，同一Block内的所有线程都能访问。延迟比Global Memory低得多（大约20-30个周期 vs 400-800个周期）。
__global__ void sharedMemKernel(float *in, float *out, int N) {
__shared__ float sdata[256]; // 声明共享内存数组
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x + tid;
// 从Global Memory加载到Shared Memory
sdata[tid] = in[i];
__syncthreads(); // 同步：确保所有线程都加载完毕
// 在Shared Memory中做计算
for (int stride = 1; stride < blockDim.x; stride *= 2) {
if (tid % (2 * stride) == 0) {
sdata[tid] += sdata[tid + stride];
}
__syncthreads();
}
// 写回结果
if (tid == 0) out[blockIdx.x] = sdata[0];
}这段代码是典型的归约（Reduction）操作。先把数据从Global Memory搬到Shared Memory，然后在Shared Memory里做累加。为什么要这么麻烦？因为如果直接在Global Memory里做，每次累加都要访问显存，慢得离谱。
核心原则：数据复用越多，Shared Memory的价值越大。如果每个数据只用一次，那Shared Memory反而成了累赘（因为多了一次拷贝）。
3.3.4 全局内存（Global Memory）
Global Memory就是显存。所有线程都能访问，但延迟最高。容量最大（几GB到几十GB），但带宽有限。
访问Global Memory时，有个重要的优化技巧：合并访问（Coalesced Access）。意思是，同一个Warp内的32个线程，如果访问的地址是连续的，硬件会把32次访问合并成一次大的内存事务，效率极高。
// 合并访问（高效）
float val = data[threadIdx.x + blockIdx.x * blockDim.x];
// 非合并访问（低效）
float val = data[threadIdx.x * 32 + blockIdx.x];第一种写法，线程0访问data[0]，线程1访问data[1]……地址连续，合并访问。第二种写法，线程0访问data[0]，线程1访问data[32]……地址间隔32，无法合并，性能差很多。
| 内存类型 | 位置 | 访问范围 | 延迟 | 容量 | 生命周期 |
|---|---|---|---|---|---|
| 寄存器 | SM内部 | 单个线程 | ~1周期 | 几十KB/SM | 线程生命周期 |
| Local Memory | 显存（有缓存） | 单个线程 | ~100周期 | 取决于显存 | 线程生命周期 |
| Shared Memory | SM内部 | 同一Block | ~20-30周期 | 几十KB/SM | Block生命周期 |
| Global Memory | 显存 | 所有线程+Host | ~400-800周期 | 几GB~几十GB | 程序生命周期 |
3.4 小结：怎么选？
说了这么多，到底怎么用？我总结几条实战经验：
- 频繁访问的数据：放寄存器。比如循环变量、临时计算结果。
- Block内需要共享的数据：放Shared Memory。比如矩阵分块计算中的子块。
- 大块数据、需要持久保存的：放Global Memory。注意合并访问。
- 尽量避免：Local Memory溢出。如果发现性能异常，用
--ptxas-options=-v编译选项查看寄存器使用情况。
嗯，这里要注意：Shared Memory和寄存器都是稀缺资源。一个SM的资源是固定的，Block越多，每个Block分到的资源就越少。有时候减少Block数量，反而能让每个Block用上更多Shared Memory，性能反而提升。这需要实际测试才能找到最优配置。
好了，这一节的内容就到这。CUDA的线程层次和内存层次是后续所有优化的基础。下一节咱们会深入Kernel的编写技巧和性能调优方法。记住我这句话：搞懂了数据怎么放、线程怎么组织，CUDA编程就学会了一半。
公众号:蓝海资料掘金营，微信deep3321