3. CUDA编程基础回顾：线程层次结构、内存层次结构、SM与Warp调度、CUDA流与事件
各位同学，咱们今天聊点实在的。CUDA编程基础，说白了就是理解GPU怎么干活。我见过太多人一上来就写kernel，结果性能惨不忍睹。为什么？因为没搞懂GPU的「脾气」。这一节，我把几个核心概念掰开揉碎了讲，都是我在实际项目中踩过的坑。
3.1 线程层次结构：从Grid到Thread
GPU的线程组织方式，有点像军队编制。Grid是集团军，Block是师，Thread就是士兵。你写代码时，得告诉GPU：我要多少个集团军，每个集团军下辖几个师，每个师有多少士兵。
我个人习惯用一维或二维的Grid和Block。三维？说实话，大部分场景用不上，反而容易把索引搞乱。
// 典型的线程层次结构定义
dim3 gridDim(16, 16); // 16x16 = 256个Block
dim3 blockDim(32, 32); // 32x32 = 1024个Thread/Block
kernel<<<gridDim, blockDim>>>(...);这里有个关键点：Block内的线程数不能超过1024。这是硬件限制，别问我为什么，老黄就是这么设计的。我在项目中遇到过有人设了2048，结果kernel直接罢工，查了半天才发现是这问题。
核心要点：线程索引计算是基本功。用 threadIdx.x、blockIdx.x、blockDim.x 这三个变量，就能算出全局唯一的线程ID。
// 全局线程ID计算（一维情况）
int tid = blockIdx.x * blockDim.x + threadIdx.x;二维或三维的情况，原理一样，就是多算几个维度。你想想看，每个线程拿到自己的ID，就知道该处理哪块数据了。这就是SIMT（单指令多线程）的精髓。
3.2 内存层次结构：速度与容量的博弈
GPU的内存层次，说白了就是「越靠近计算单元越快，但容量越小」。我刚开始做算子移植时，总想把所有数据都塞进共享内存，结果发现根本塞不下。后来学乖了，得根据数据访问模式来选。
| 内存类型 | 位置 | 速度 | 容量 | 作用域 |
|---|---|---|---|---|
| 全局内存 | 显存 | 慢（~200周期） | 大（GB级） | 所有线程 |
| 共享内存 | SM内部 | 快（~5周期） | 小（48KB~164KB） | Block内 |
| 寄存器 | SM内部 | 极快（1周期） | 极小（每个线程最多255个） | 单个线程 |
| 常量内存 | 显存（带缓存） | 快（有缓存） | 小（64KB） | 所有线程（只读） |
嗯，这里要注意：共享内存是手动管理的。你得自己把数据从全局内存搬到共享内存，用完了还得考虑同步。我曾经在移植一个矩阵乘法算子时，忘了加 __syncthreads()，结果数据竞争导致结果全错。那叫一个惨。
我的经验：如果数据被同一个Block内的多个线程反复访问，果断用共享内存。如果每个线程只访问自己的数据，寄存器就够了。别什么都往共享内存塞，容量有限，而且搬数据也有开销。
3.3 SM与Warp调度：理解GPU的「小分队」
SM（Streaming Multiprocessor）是GPU的核心计算单元。每个SM里有多个CUDA Core，还有共享内存、寄存器文件、调度器等。你写的Block会被分配到SM上执行，一个SM可以同时跑多个Block。
但真正执行指令的单位是Warp。Warp是32个线程组成的小分队，它们执行相同的指令。说白了，GPU不是一次调度一个线程，而是一次调度一个Warp。
关键洞察：如果Warp内的32个线程走不同的分支（比如if-else），那这个Warp就得串行执行所有分支。这就是所谓的「Warp Divergence」。性能杀手，没有之一。
我举个例子。假设你写了个kernel，里面有个if语句：
if (tid % 2 == 0) {
// 偶数线程执行这里
} else {
// 奇数线程执行这里
}同一个Warp里，一半线程走if，一半走else。GPU会先执行if分支，把else的线程屏蔽掉；再执行else分支，把if的线程屏蔽掉。效率直接砍半。我在项目中遇到过这种情况，优化前性能只有理论值的60%。后来把分支条件改成按Warp对齐，性能直接翻倍。
那怎么避免？尽量让Warp内的线程走相同的路径。如果必须用分支，让分支条件基于 warpId = threadIdx.x / 32 而不是 threadIdx.x % 2。
3.4 CUDA流与事件：并发执行的艺术
CUDA流，说白了就是任务队列。你把不同的操作放到不同的流里，GPU可以并行执行它们。比如一边做计算，一边做数据传输。
我刚开始用流时，以为只要创建多个流就能自动并行。结果发现，如果流之间没有依赖关系，确实可以并行；但如果一个流要等另一个流的结果，那就得用事件来同步。
cudaStream_t stream1, stream2;
cudaEvent_t event;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaEventCreate(&event);
// 在stream1中执行kernel
kernel1<<<grid, block, 0, stream1>>>(...);
// 记录事件
cudaEventRecord(event, stream1);
// 在stream2中等待事件，然后执行kernel2
cudaStreamWaitEvent(stream2, event);
kernel2<<<grid, block, 0, stream2>>>(...);这里有个坑：事件记录和等待必须在正确的流上。我曾经把事件记录在stream1，却在stream3上等待，结果stream3根本没等到，数据就乱了。嗯，调试这种问题特别痛苦，因为错误不是每次都复现。
避坑指南：我曾经在移植一个多流算子时，忘了检查流是否已经完成就释放了流资源。结果程序崩溃，查了两天才发现是资源释放顺序的问题。记住：cudaStreamDestroy 会阻塞直到流中所有操作完成，但如果你在另一个线程里操作同一个流，就可能出问题。
事件还有一个妙用：测量时间。用 cudaEventElapsedTime 可以精确测量kernel的执行时间，比用CPU计时器准多了。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
kernel<<<grid, block>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel耗时: %f ms\n", ms);3.5 知识体系总览
说了这么多，咱们用一张图把整个知识体系串起来。这张图是我自己画的，涵盖了CUDA编程的核心逻辑。
这张图把四个核心概念串起来了。从上到下，线程层次决定了「谁干活」，内存层次决定了「数据放哪」，SM与Warp决定了「怎么调度」，流与事件决定了「怎么并发」。你写任何CUDA程序，都绕不开这四个维度。
好了，这一节的内容就到这。记住：理论是基础，但真正的高手都是在实践中磨出来的。多写代码，多踩坑，自然就懂了。