GPU架构基础：SIMT模型、Warp调度、内存层次结构
好，咱们正式开始聊GPU架构。说实话，很多搞图形或者搞计算的同学，一开始就把GPU当成一个「黑盒子」——数据扔进去，结果出来，中间发生了什么？不清楚。我个人习惯是，不管用啥工具，先把它的「脾气」摸透。GPU的脾气，说白了就三个关键词：SIMT、Warp、内存层次。搞懂这三个，你写出来的代码效率能差出好几倍。
一、SIMT模型：不是你想的那种SIMD
很多人一上来就说「GPU是SIMD架构」。嗯，这话对了一半。GPU用的是SIMT（Single Instruction, Multiple Threads），跟CPU上的SIMD有本质区别。
CPU的SIMD，比如AVX指令集，是你写一条指令，它同时对一批数据做同样的操作。但GPU的SIMT呢？它是一条指令，驱动多个独立的线程去执行。每个线程有自己的寄存器、自己的程序计数器。你想想看，这就像教官喊一声「齐步走」，每个士兵虽然步伐一致，但每个人都有自己的想法（数据）。
核心区别：
- SIMD：单指令，单线程，多数据。程序员手动控制向量化。
- SIMT：单指令，多线程，多数据。硬件自动管理线程并行。
我在项目中遇到过一个问题：有个同事把CPU上的SIMD优化思路直接搬到GPU上，结果性能反而下降了。为什么？因为SIMT的线程是「逻辑上并行，物理上串行」的。你写32个线程，硬件可能只用一个CUDA Core去轮流执行。嗯，这里要注意，SIMT的「线程」是轻量级的，创建和切换开销极低，但你不能假设它们真的在同时跑。
二、Warp调度：GPU的「最小执行单元」
好，线程有了，谁来管它们？答案是Warp。
在NVIDIA的GPU里，32个线程组成一个Warp。这是硬件调度的基本单位。什么意思呢？就是SM（Streaming Multiprocessor）每次取一条指令，然后让一个Warp里的32个线程同时执行。如果这32个线程走的是同一个分支，那效率最高。如果它们走了不同的分支，那就麻烦了——这叫Warp Divergence。
避坑指南：我曾经在一个图像处理项目里，写了个if-else判断像素类型。结果一个Warp里的32个像素，16个走if，16个走else。硬件只能先执行if的那16个，再执行else的那16个。性能直接砍半。后来我改成用? :三元运算符，或者用__ballot_sync做谓词判断，才把性能拉回来。
Warp调度还有个特点：零开销线程切换。CPU切换线程要保存上下文，GPU不用。因为每个Warp的寄存器都是物理上独立的。SM里同时驻留几十个Warp，当一个Warp在等内存访问结果时，硬件立刻切换到另一个Warp。这就是GPU能隐藏延迟的秘诀。
我个人习惯，写CUDA代码时，会先算一下每个SM能驻留多少个Warp。这个数叫Occupancy。Occupancy越高，延迟隐藏能力越强。但也不是越高越好，有时候为了给每个线程多分配寄存器，得牺牲一些Occupancy。这是个trade-off。
三、内存层次结构：从快到慢，从贵到便宜
GPU的内存，跟CPU一样，也是个金字塔结构。但GPU的层次更复杂，也更讲究「因地制宜」。
| 内存类型 | 位置 | 访问延迟 | 作用域 | 典型用途 |
|---|---|---|---|---|
| 寄存器 | SM内部 | 1 cycle | 单个线程 | 局部变量、中间计算结果 |
| 共享内存 | SM内部 | ~5 cycles | Block内所有线程 | 线程间数据共享、归约操作 |
| L1缓存 | SM内部 | ~10 cycles | SM内所有线程 | 全局内存的缓存 |
| 全局内存 | 显存 | ~400 cycles | 所有线程 | 主数据存储 |
| 常量内存 | 显存（带缓存） | ~10 cycles（缓存命中） | 所有线程 | 只读常量数据 |
| 纹理内存 | 显存（带缓存） | ~100 cycles | 所有线程 | 空间局部性强的数据 |
你看这个表，从寄存器到全局内存，延迟差了400倍。你想想看，如果代码里频繁访问全局内存，那性能肯定好不了。
我建议，写CUDA Kernel时，先想清楚数据流：
- 全局内存 → 共享内存：把要用的数据一次性搬进共享内存。
- 共享内存 → 寄存器：在共享内存里做计算，结果放寄存器。
- 寄存器 → 全局内存：计算完再写回全局内存。
这个流程，说白了就是「数据本地化」。我在做矩阵乘法优化时，就是靠这个思路，把性能从几十GFLOPS提到了接近理论峰值。
四、内存访问模式：合并访问
说到全局内存，有个概念必须提——合并访问（Coalesced Access）。
GPU访问全局内存时，是以128字节为单位的。如果一个Warp里的32个线程，访问的地址是连续的，那硬件一次就能把数据取回来。如果地址是分散的，那就要多次访问，性能就差了。
小技巧：写代码时，尽量让线程ID和内存地址对应起来。比如data[threadIdx.x]就是合并访问，而data[threadIdx.x * 32]就是非合并的。我习惯在Kernel开头加个断言，检查地址对齐情况。
举个例子：
// 合并访问（推荐）
float val = data[threadIdx.x + blockIdx.x * blockDim.x];
// 非合并访问（避免）
float val = data[threadIdx.x * 32 + blockIdx.x];你看，第一个写法，线程0访问地址0，线程1访问地址1，连续。第二个写法，线程0访问地址0，线程1访问地址32，中间隔了31个元素。硬件得拆成32次访问，性能直接掉到1/32。
五、总结一下
嗯，这一章的内容其实就三句话：
- SIMT模型：别把它当SIMD用，每个线程是独立的。
- Warp调度：避免分支发散，提高Occupancy。
- 内存层次：多用共享内存，少用全局内存，注意合并访问。
我记得刚学CUDA那会儿，觉得这些概念太抽象。后来在项目里踩了坑，才真正理解。你如果现在觉得有点懵，没关系，后面几章我们会结合具体案例，把这些概念一个一个用起来。到时候你就知道，这些「基础」才是真正决定性能上限的东西。
下一章，我们聊聊CUDA编程模型——怎么把CPU上的串行代码，改造成GPU上的并行Kernel。到时候我会拿一个实际项目里的例子，手把手带你走一遍。