第4章：CUDA编程基础回顾：线程层次结构、内存层次结构、SM与warp调度
好，咱们今天聊聊CUDA编程里最核心的几个概念。说实话，很多同学一开始学CUDA，就被各种线程、块、网格搞得晕头转向。我当年也一样，第一次写kernel的时候，连线程ID怎么算都搞错了，结果跑出来一堆乱码。嗯，今天咱们就把这些东西彻底理清楚。
4.1 线程层次结构：从Grid到Thread
CUDA的线程模型，说白了就是三层结构：Grid（网格）→ Block（线程块）→ Thread（线程）。你想想看，这就像是一个军队的编制：Grid是军团，Block是连队，Thread就是每个士兵。
关键点：每个线程都有唯一的ID，通过 threadIdx、blockIdx、blockDim 和 gridDim 来定位。
我个人习惯用一维索引来简化问题，但实际项目中经常用到二维甚至三维。比如处理图像时，用二维线程块就特别自然——每个线程对应一个像素。
// 一个典型的kernel启动方式
dim3 gridDim(16, 16); // 16x16 个block
dim3 blockDim(32, 32); // 每个block有32x32个线程
myKernel<<<gridDim, blockDim>>>(d_data, width, height);
// 在kernel内部计算全局线程ID
int tx = threadIdx.x + blockIdx.x * blockDim.x;
int ty = threadIdx.y + blockIdx.y * blockDim.y;
int globalIdx = ty * width + tx;小技巧：blockDim的大小通常选32的倍数，因为warp大小就是32。选64或128比较常见，太大反而可能降低occupancy。
4.2 内存层次结构：速度与容量的博弈
CUDA的内存模型，其实就是一个金字塔结构。越往上速度越快，但容量越小。我刚开始做优化时，总想着把所有数据都塞进共享内存，结果发现根本放不下——这就是没理解这个层次结构。
| 内存类型 | 位置 | 访问延迟 | 作用域 | 生命周期 |
|---|---|---|---|---|
| 寄存器 | SM内部 | ~1 cycle | 单个线程 | kernel执行期间 |
| 共享内存 | SM内部 | ~5 cycles | 单个block | block生命周期 |
| 全局内存 | 显存 | ~400 cycles | 所有线程 | 程序运行期间 |
| 常量内存 | 显存（带缓存） | ~10 cycles（缓存命中） | 所有线程 | 程序运行期间 |
| 纹理内存 | 显存（带缓存） | ~10 cycles（缓存命中） | 所有线程 | 程序运行期间 |
这里我要特别强调一下共享内存。它虽然只有几十KB，但用好了能带来几十倍的性能提升。我曾经在优化一个矩阵转置的算子时，就是因为合理使用了共享内存做数据分块，把带宽利用率从30%提到了85%以上。
注意：共享内存有bank conflict的问题。如果多个线程访问同一个bank的不同地址，就会产生冲突，导致串行化。解决办法通常是加padding或者调整访问模式。
4.3 SM与warp调度：硬件怎么执行你的代码
SM（Streaming Multiprocessor）是GPU的核心计算单元。每个SM里有多个CUDA Core，还有共享内存、寄存器文件、warp调度器等。你写的kernel，最终就是在这些SM上执行的。
warp是SM调度的基本单位，大小固定为32个线程。同一个warp里的线程执行相同的指令（SIMT模型）。但要注意，如果线程有分支，就会产生warp divergence——说白了就是有些线程走if，有些走else，导致部分线程被屏蔽，性能下降。
我记得有一次调试一个复杂的计算图，发现某个算子的性能死活上不去。后来用nvprof一看，warp divergence率高达40%。原来是条件判断里有个变量是线程ID相关的，导致每个warp里都有分支。改成用位运算统一处理之后，性能直接翻倍。
Occupancy（占用率）：每个SM能同时容纳的warp数量与最大warp数量的比值。occupancy越高，越能隐藏内存访问延迟。但也不是越高越好——有时候降低occupancy换取更多寄存器，反而能提升单线程性能。
下面这张图展示了SM内部的结构和warp调度流程：
这张图展示了SM内部的核心组件。warp调度器会从多个就绪的warp中选择一个，然后通过dispatch unit把指令发射到CUDA Core上执行。共享内存和寄存器文件则提供快速的数据访问。
4.4 实战中的避坑指南
说了这么多理论，咱们来点实际的。我在优化算子时总结了几条经验：
- 线程块大小选256或512：太小了SM利用率低，太大了寄存器压力大。256和512是大多数算子的甜点值。
- 尽量减少全局内存访问：能放共享内存就放共享内存，能合并访问就合并访问。我曾经遇到一个算子，80%的时间都花在全局内存访问上，改成共享内存分块后快了5倍。
- 注意warp divergence：如果条件判断依赖于线程ID，尽量用位运算或三元运算符替代if-else。
- 用
__syncthreads()同步：在共享内存读写之间一定要加同步，否则会出现数据竞争。我踩过这个坑，debug了整整两天才发现是少了个同步。
个人经验：刚开始做CUDA优化时，别急着调occupancy。先保证功能正确，然后用nvprof/ncu看瓶颈在哪里。很多时候，内存访问模式比occupancy更重要。
好了，这一章的内容就到这里。线程层次结构、内存层次结构、SM与warp调度，这三个概念是CUDA编程的基石。理解透了，后面的算子优化才能事半功倍。
一句话总结：线程层次结构决定了你怎么组织计算，内存层次结构决定了数据怎么流动，SM和warp调度决定了硬件怎么执行。三者协同，才能写出高效的CUDA程序。
想要更多资料，联系可可