1. CUDA编程模型回顾：从Grid到Thread，你真正理解了吗？
说实话，很多同学学CUDA时，一上来就写kernel函数，结果性能惨不忍睹。我当年也是这样——代码能跑，但就是慢。后来才发现，问题出在最基础的编程模型理解上。今天咱们就把这块地基打牢。
1.1 Grid-Block-Thread：GPU的“三级组织架构”
想象一下，你要处理一张4K图片，每个像素都要做计算。CPU的做法是：一个核心吭哧吭哧算完所有像素。GPU呢？它把任务拆成三级：
- Grid（网格）：整个任务，比如整张图片
- Block（线程块）：任务的一个子集，比如图片的一块区域
- Thread（线程）：最小的执行单元，处理一个像素
我习惯用军队来类比：Grid是集团军，Block是连队，Thread就是每个士兵。指令从统帅部（CPU）下达，但具体怎么打，每个连队内部可以灵活配合。
关键点：Block内的线程可以协作（通过共享内存），但不同Block的线程是“老死不相往来”的。这是硬件决定的，别想着跨Block通信——代价极高。
来看一个典型的kernel启动代码：
// 定义网格和块维度
dim3 gridDim(32, 32); // 32x32 = 1024个block
dim3 blockDim(16, 16); // 每个block 16x16 = 256个线程
// 启动kernel
myKernel<<<gridDim, blockDim>>>(d_data, width, height);这里有个坑：blockDim的乘积不能超过1024（对于大多数GPU）。我刚开始写代码时，设了个32x32=1024的block，刚好卡在极限上。后来换到老一点的GPU，直接报错。嗯，这里要注意——不同架构的极限值不一样，查一下你的GPU的compute capability。
1.2 内存层次结构：谁快谁慢，心里要有数
GPU的内存就像一座金字塔。塔尖最快，但容量最小；塔基最慢，但容量最大。咱们从快到慢捋一遍：
| 内存类型 | 位置 | 速度 | 容量 | 作用域 |
|---|---|---|---|---|
| 寄存器 (Register) | SM内部 | ~1周期 | 每个线程最多255个 | 单个线程 |
| 共享内存 (Shared Memory) | SM内部 | ~5周期 | 48KB~164KB | 单个Block |
| 全局内存 (Global Memory) | 显存 | ~400周期 | 数GB | 所有线程 |
| 局部内存 (Local Memory) | 显存（但被编译器隐藏） | ~400周期 | 有限 | 单个线程 |
你想想看，全局内存访问一次要400个周期，而寄存器只要1个周期。这差距有多大？说白了，优化CUDA程序的核心，就是尽量减少全局内存访问，把数据搬到共享内存或寄存器里算。
个人经验：我做过一个矩阵乘法的优化项目，一开始直接读写全局内存，性能只有20 GFLOPS。后来用了共享内存做分块（tiling），直接飙到800 GFLOPS。40倍的提升，就靠这一招。
关于局部内存，很多新手会忽略。其实它是个“隐形杀手”——当你声明一个数组，但编译器发现寄存器不够用时，就会悄悄把它放到局部内存里。局部内存物理上在显存，速度跟全局内存一样慢。我曾经排查过一个性能问题，发现就是某个临时数组被“溢出”到了局部内存，改小数组大小后性能翻倍。
1.3 Warp执行模型：SIMT架构的精髓
这是CUDA最核心、也最容易踩坑的地方。先看一张图：
这张图展示了Warp的执行逻辑。简单说：32个线程组成一个Warp，它们共享一个指令计数器。也就是说，这32个线程在同一时刻执行同一条指令，只是操作的数据不同。
为什么会这样？因为GPU的设计哲学是“用更多的计算单元来换取并行度”，而不是像CPU那样做复杂的分支预测。所以它干脆让32个线程绑在一起，步调一致地执行。
避坑指南：我曾经写过一个kernel，里面有个if-else分支：
if (threadIdx.x % 2 == 0) {
// 偶数线程执行A
do_something_A();
} else {
// 奇数线程执行B
do_something_B();
}结果性能惨不忍睹。为什么？因为同一个Warp里的线程，有的走if，有的走else，这就叫Warp divergence（分支发散）。GPU的做法是：先执行所有走if的线程（另一半闲着），再执行所有走else的线程（另一半闲着）。相当于性能直接砍半！
怎么避免？我总结了几个原则：
- 尽量让Warp内所有线程走同一个分支。比如用
threadIdx.x < 16做条件，前16个线程走if，后16个走else——这没问题，因为它们在同一个Warp里是连续的。 - 如果必须分支，考虑用三元运算符或位运算代替。比如
result = (condition) ? val_a : val_b，编译器有时能优化成无分支代码。 - 实在不行，就重新组织数据。把需要走同一分支的数据打包在一起处理。
1.4 内存合并访问：让全局内存“飞”起来
说到全局内存，还有个重要概念——合并访问（coalesced access）。说白了，就是让同一个Warp里的32个线程，访问连续的内存地址。这样GPU可以把这些访问合并成一次大的内存事务，效率极高。
来看两个例子：
// 好的：合并访问
float val = data[threadIdx.x]; // 线程0访问data[0]，线程1访问data[1]...
// 坏的：非合并访问
float val = data[threadIdx.x * 32]; // 线程0访问data[0]，线程1访问data[32]...第一个例子，32个线程访问连续地址，一次内存事务搞定。第二个例子，地址间隔32个元素，需要多次事务，性能差很多。
我的习惯：写kernel时，我会先检查内存访问模式。如果发现非合并访问，就考虑转置数据或者调整线程索引的计算方式。这一步优化往往能带来2-3倍的性能提升。
1.5 小结：把基础打牢，后面才走得远
好了，这一章的内容就这些。咱们回顾一下核心要点：
- Grid-Block-Thread是三级并行模型，Block内线程可以协作，Block间不行
- 内存层次从快到慢：寄存器 → 共享内存 → 全局内存，优化目标是把数据往快了搬
- Warp是32个线程的捆绑执行单元，避免分支发散是关键
- 全局内存访问要尽量合并，让Warp内的线程访问连续地址
这些概念看起来简单，但真正用好需要大量实践。我刚开始学的时候，也是反复踩坑才慢慢理解的。别急，后面几章咱们会深入每个优化技巧，到时候你会对这些基础有更深的体会。
想要更多资料，联系可可