第一章：CUDA编程模型回顾——从Grid到Warp的完整脉络
做GPU优化这么多年，我见过太多人一上来就盯着代码调优，结果连最基本的线程模型都没搞明白。说实话，CUDA编程模型就像盖房子的地基——地基不稳，后面再怎么装修也是白搭。今天咱们就把这块地基彻底夯实了。
1.1 Grid-Block-Thread层级结构
先看这张图，它概括了CUDA最核心的层级关系：
这张图我画了很多遍，每次给团队新人讲课时都会拿出来。你看，最外层是Grid，里面是Block，最里面是Thread。这个三层结构，说白了就是GPU用来组织并行任务的骨架。
Grid 是整个任务的总称。你调用一个kernel函数，就启动了一个Grid。Grid可以是一维、二维或三维的，但实际项目中我99%的情况只用一维或二维——三维用得少，除非处理体数据。
Block 是中间层，也是优化的关键。每个Block内的线程可以协作——通过共享内存和同步操作。我记得刚入行时犯过一个错误：把Block设得太大，结果共享内存不够用，程序直接崩了。后来养成了习惯，先算好共享内存预算再定Block大小。
Thread 是最小执行单元。每个线程执行相同的kernel代码，但处理不同的数据——这就是SIMT的精髓。你想想看，成千上万个线程同时跑，每个线程拿自己的索引去取数据，效率能不高吗？
关键约束：
- 每个Block最多1024个线程（现代GPU，老架构是512或768）
- Grid维度最大为2^31-1（一维）或65535×65535（二维）
- Block内的线程通过
__syncthreads()同步
1.2 内存层次结构——谁快谁慢，心里要有数
做GPU优化，说白了就是在跟内存打交道。我见过太多人把数据一股脑扔进全局内存，然后抱怨GPU慢。其实GPU的内存层次很丰富，用对了地方，性能能翻好几倍。
| 内存类型 | 位置 | 访问延迟 | 生命周期 | 作用域 |
|---|---|---|---|---|
| Register | SM内部 | ~1 cycle | 线程内 | 单个线程 |
| Shared Memory | SM内部 | ~5 cycles | Block内 | Block内所有线程 |
| Local Memory | 显存（DRAM） | ~100 cycles | 线程内 | 单个线程（寄存器溢出时） |
| Global Memory | 显存（DRAM） | ~400-800 cycles | 整个程序 | 所有线程 + Host |
| Constant Memory | 显存（带缓存） | ~100 cycles（缓存命中时更快） | 整个程序 | 所有线程（只读） |
| Texture Memory | 显存（带缓存） | ~100 cycles | 整个程序 | 所有线程（只读，支持插值） |
这张表我建议你打印出来贴在工位上。每次写kernel之前看一眼，问自己三个问题：
- 这个数据能不能放寄存器？能就别放共享内存。
- 这个数据能不能放共享内存？能就别放全局内存。
- 这个数据是不是只读且所有线程都访问？是就考虑常量内存。
我的经验： 寄存器是最快的，但数量有限。每个线程最多255个寄存器（现代GPU）。如果你写的代码用了太多局部变量，编译器会把多余的变量"溢出"到Local Memory——这玩意儿其实在显存里，慢得很。我曾经优化过一个矩阵乘法的kernel，就是靠减少寄存器使用量，把性能从60%提升到了85%。
1.3 Shared Memory——Block内的"小团队协作区"
共享内存是CUDA优化的核心武器。它位于SM内部，速度比全局内存快两个数量级。但它的容量很小——通常每个SM只有48KB到164KB（取决于架构）。
共享内存的典型用法：
- 数据复用： 多个线程需要访问同一块数据时，先加载到共享内存
- 合并访问： 把非合并的全局内存访问模式，转成合并的共享内存访问
- 线程间通信： Block内的线程通过共享内存交换数据
// 共享内存的典型用法：矩阵分块
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int row = by * TILE_SIZE + ty;
int col = bx * TILE_SIZE + tx;
float sum = 0.0f;
for (int t = 0; t < N / TILE_SIZE; t++) {
// 协作加载数据到共享内存
As[ty][tx] = A[row * N + t * TILE_SIZE + tx];
Bs[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];
__syncthreads();
// 从共享内存计算
for (int k = 0; k < TILE_SIZE; k++) {
sum += As[ty][k] * Bs[k][tx];
}
__syncthreads();
}
C[row * N + col] = sum;
}注意： 共享内存使用不当会导致 bank conflict。共享内存被分成32个bank（每个bank 4字节宽），如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。我踩过这个坑——一个简单的转置kernel，因为没处理好bank conflict，性能直接腰斩。
1.4 Warp执行模型与SIMT架构
这是CUDA最核心、也最容易让人困惑的概念。说白了，GPU并不是真的同时执行所有线程——它把32个线程打包成一个Warp，以Warp为单位调度执行。
为什么是32？ 这是NVIDIA的设计选择。32个线程共享一个指令单元，执行相同的指令。如果线程之间分支不同（比如if-else），就会发生Warp Divergence——一部分线程执行，另一部分等待，性能直接打折扣。
Warp Divergence 示例：
// 不好的写法——会导致Warp Divergence
if (threadIdx.x % 2 == 0) {
// 偶数线程执行这里
result = a + b;
} else {
// 奇数线程执行这里
result = a * b;
}
// 好的写法——避免分支
int mask = threadIdx.x & 1; // 0或1
result = mask ? (a * b) : (a + b);我曾经在一个图像处理项目里，因为一个简单的条件判断，导致Warp内一半线程闲置，性能掉了40%。后来改成用位运算和三元表达式，问题就解决了。你想想看，32个线程里如果有16个走分支A、16个走分支B，那这16个等待的时候，计算单元就空转了。
SIMT vs SIMD：
- SIMD（单指令多数据）： 所有处理单元必须执行相同指令，数据宽度固定（如SSE的128位）
- SIMT（单指令多线程）： 每个线程有自己的程序计数器，可以独立分支（但分支会导致性能损失）
SIMT更灵活，但代价就是Warp Divergence。所以写CUDA代码时，要尽量让同一个Warp内的线程走相同的执行路径。
1.5 实战中的避坑指南
说了这么多理论，最后分享几个我在项目中踩过的坑：
避坑1： 我曾经以为Block越大越好，结果把Block设成了1024，但每个线程用了太多寄存器，导致Occupancy（占用率）只有25%。后来把Block改成256，Occupancy提升到75%，性能反而翻倍了。记住：不是线程越多越好，是活跃的Warp越多越好。
避坑2： 共享内存的bank conflict不是总能避免的。如果实在避不开，可以尝试padding——在共享内存数组的每一行末尾加几个空元素，改变地址对齐方式。我有个同事靠这招把性能提升了30%。
避坑3： 全局内存的合并访问。Warp访问全局内存时，如果32个线程访问的地址是连续的（且对齐到128字节），硬件会合并成一次或几次内存事务。否则，每个线程单独访问，效率极低。我见过有人写 array[threadIdx.x * 32] 这种模式，结果性能惨不忍睹。
好了，这一章的内容就到这里。记住这些基础概念，后面我们讲具体优化技巧时，你会发现自己能更快地理解问题出在哪里。