GPU架构基础：CUDA核心概念、SM与Warp调度、内存层次结构
各位同学，今天我们来聊聊GPU架构。说实话，我刚开始接触CUDA时，也被那一堆术语搞得头晕——SM是什么？Warp又是什么？Global Memory和Shared Memory到底差在哪？
别急，我们一个一个来拆解。搞懂了这些，你写算子的时候就能知道瓶颈在哪，优化方向在哪。
一、CUDA核心概念：从线程到网格
CUDA的编程模型，说白了就是「分层」的。我习惯把它想象成一个军队的编制：
- Thread（线程）：最基层的士兵，每个线程执行一条指令
- Block（线程块）：一个班，里面的线程可以互相通信（通过Shared Memory）
- Grid（网格）：一个连队，由多个Block组成
你写一个kernel函数，实际上是在描述「每个线程该干什么」。GPU会帮你把这些线程分配到硬件上执行。
关键点：CUDA的编程模型是逻辑上的，而硬件执行是物理上的。你写的Grid/Block只是告诉编译器「我有这么多工作要做」，至于怎么分给SM，那是硬件的事。
举个例子，你要做1000个元素的向量加法：
__global__ void vec_add(float *a, float *b, float *c, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
// 启动配置：256个线程一个Block，一共 ceil(1000/256) = 4个Block
vec_add<<<4, 256>>>(a, b, c, 1000);这里每个线程处理一个元素。嗯，简单粗暴，但很有效。
二、SM与Warp调度：GPU的「心脏」
SM（Streaming Multiprocessor）是GPU真正的计算单元。你可以把SM想象成一个「小型的多核CPU」，但它有自己的调度逻辑。
Warp是什么？
Warp是SM调度的基本单位，通常包含32个线程。为什么是32？这是NVIDIA的设计选择，跟指令宽度、寄存器文件有关。
我个人习惯把Warp理解成「一捆线程」——这32个线程执行同一条指令，只是操作不同的数据。这就是SIMT（单指令多线程）的精髓。
避坑指南：我曾经在写一个reduce算子时，没注意Warp内线程的分支发散。结果32个线程里有一半走if，一半走else，性能直接腰斩。记住：同一个Warp里的线程，尽量走同一个分支。
SM怎么调度Warp？
SM里有一个Warp Scheduler，它维护着一堆Warp的状态。每个时钟周期，Scheduler会选一个「就绪」的Warp发射指令。
为什么会这样？因为内存访问有延迟。当一个Warp在等数据时，Scheduler就切换到另一个Warp去执行。这叫「零开销上下文切换」——切换Warp不需要保存恢复寄存器，因为每个Warp有自己的寄存器文件。
你想想看，这跟CPU的线程切换完全不一样。CPU切换线程要保存一堆上下文，GPU不用。所以GPU可以同时维护很多Warp来隐藏延迟。
下面我用一张图来展示SM的内部结构：
这张图展示了SM内部的核心组件。Warp Scheduler从Warp Pool里选一个，然后从指令缓存取指令，发给执行单元去算。
三、内存层次结构：速度与容量的博弈
GPU的内存跟CPU一样，也有层次结构。但GPU的层次更「极端」——离计算单元越近越快，但容量越小。
| 内存类型 | 位置 | 大小 | 延迟 | 作用域 |
|---|---|---|---|---|
| Register | SM内部 | 每个SM约64K~256K个 | ~1 cycle | 单个线程 |
| Shared Memory | SM内部 | 每个SM 48KB~164KB | ~5 cycles | 单个Block |
| Global Memory | 显存（DRAM） | 几GB~几十GB | ~400 cycles | 所有线程 |
看到这个延迟差距了吗？Register只要1个时钟周期，Global Memory要400个。差了400倍！
Register（寄存器）
寄存器是每个线程私有的。你在kernel里定义的局部变量，只要不溢出，都会放在寄存器里。
注意：寄存器是有限的！每个SM的寄存器总数是固定的。如果你一个Block用了太多寄存器，SM能同时容纳的Block数就会减少，导致Occupancy下降。
我遇到过一个问题：一个kernel用了很多局部变量，结果编译器报了「spill to local memory」的警告。这意味着寄存器不够用，变量被挤到Global Memory里去了，性能直接崩了。
Shared Memory（共享内存）
Shared Memory是同一个Block内所有线程共享的。它的速度比Global Memory快得多，但需要手动管理。
说白了，Shared Memory就是「程序员可控的缓存」。你手动把数据从Global Memory搬到Shared Memory，然后让Block内的线程一起用。
举个例子，矩阵乘法中常用的tiling技术：
__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int row = blockIdx.y * TILE_SIZE + threadIdx.y;
int col = blockIdx.x * TILE_SIZE + threadIdx.x;
float sum = 0.0f;
for (int t = 0; t < N / TILE_SIZE; t++) {
// 协作加载数据到Shared Memory
As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
__syncthreads(); // 同步：确保所有线程都加载完了
// 从Shared Memory计算
for (int k = 0; k < TILE_SIZE; k++) {
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
__syncthreads();
}
C[row * N + col] = sum;
}这里的关键是__syncthreads()——它保证所有线程都完成加载后，才开始计算。不然有的线程还在等数据，有的已经开始算了，结果就错了。
避坑指南：我曾经在Shared Memory里写了一个bank conflict的代码，性能比不用Shared Memory还差。Shared Memory被分成32个bank，如果多个线程同时访问同一个bank的不同地址，就会冲突。解决办法是给数组加padding，或者调整访问模式。
Global Memory（全局内存）
Global Memory就是显存，容量大但延迟高。访问Global Memory时，要尽量做到「合并访问」——相邻的线程访问相邻的地址。
为什么？因为Global Memory的访问是以128字节为单位的。如果32个线程访问连续的128字节，一次就能搞定。如果访问分散的地址，就要多次访问，带宽就浪费了。
我刚开始写CUDA时，没注意这个，写了个随机访问的代码，结果带宽利用率只有10%。后来改成连续访问，性能直接提升了8倍。
四、总结：一张图看懂GPU架构
最后，我用一张图总结整个GPU架构的层次关系：
从Grid到Thread，每一层都有对应的硬件映射。搞懂了这张图，你就知道你的代码在GPU上是怎么跑的。
嗯，今天就先聊到这。记住：写算子时，多想想数据在哪、怎么访问、线程怎么协作。这些基础打牢了，后面的优化技巧才能用得上。
核心要点回顾：
- CUDA编程模型：Grid → Block → Thread，逻辑分层
- SM是物理执行单元，Warp是调度单位（32线程一组）
- 内存速度：Register > Shared Memory > Global Memory
- 优化方向：减少Global访问、合并访问、避免Warp发散、控制寄存器使用
想要更多资料，联系可可