3. GPU编程模型：从SIMT到内存访问的艺术
各位同学，今天我们来聊聊GPU编程模型的核心。说实话，这个主题我讲了不下五十次，但每次都有新感悟。GPU编程不像CPU编程那样直来直去，它有自己的脾气和性格。理解这些，你才能写出真正高效的GPU代码。
3.1 SIMT执行模型：看似简单，实则精妙
SIMT，全称是Single Instruction Multiple Threads。翻译过来就是「单指令多线程」。你可能会问：这和CPU的SIMD有什么区别？
嗯，区别大了。SIMD是数据并行，一条指令处理多个数据。而SIMT是线程并行，一条指令驱动多个线程。每个线程有自己的程序计数器、寄存器和执行状态。说白了，SIMT是硬件帮你做了线程调度，你只管写代码就行。
核心要点：SIMT模型中，32个线程组成一个warp（在NVIDIA术语中）。这32个线程执行同一条指令，但处理不同的数据。如果遇到分支，warp会串行执行所有分支路径。
我在项目中遇到过一个问题：一个简单的if-else语句，让性能直接腰斩。为什么？因为warp中的线程走了不同的分支，硬件不得不分别执行两条路径。这就是所谓的「warp divergence」。你想想看，32个线程本来可以一起干活，现在却要分成两拨，效率自然就下来了。
避坑指南：我曾经在图像处理项目中，因为一个边界检查的if语句，导致性能损失了40%。后来我把边界检查挪到了循环外面，用掩码操作替代分支，性能立刻上来了。记住：尽量让warp内的线程走相同的执行路径。
3.2 线程层次结构：Grid、Block、Thread
GPU的线程组织方式，说白了就是三级结构：Grid → Block → Thread。这就像军队的编制：军、师、班。每个层级都有自己的职责和限制。
| 层级 | 说明 | 典型限制 |
|---|---|---|
| Grid | 整个任务，包含多个Block | 最大维度：2^31-1（一维） |
| Block | 线程组，共享共享内存 | 最大线程数：1024（现代GPU） |
| Thread | 最小执行单元 | 每个线程有自己的寄存器和局部内存 |
我个人习惯这样理解：Grid对应整个数据集，Block对应数据分块，Thread对应单个数据元素。举个例子，处理一张1920×1080的图片，你可以用32×32的Block，每个线程处理一个像素。
// 典型的CUDA核函数调用
dim3 gridDim(64, 34); // 64×34 = 2176个Block
dim3 blockDim(32, 32); // 32×32 = 1024个线程/Block
kernel<<<gridDim, blockDim>>>(input, output, width, height);这里要注意：Block内的线程可以通过共享内存通信，但Block之间是独立的。我刚开始学的时候总想跨Block通信，结果发现根本行不通。后来才明白，GPU的设计哲学就是「分而治之」。
3.3 内存访问模式：合并访问是关键
说到内存访问，这可是GPU编程的重头戏。GPU的内存带宽很大，但如果你访问模式不对，带宽利用率可能连10%都不到。
合并访问（Coalesced Access）是什么意思？说白了，就是让warp中32个线程访问连续的内存地址。这样硬件可以把32次内存访问合并成一次大的内存事务。你想想看，一次搞定和分32次搞定，效率能一样吗？
黄金法则：warp中线程的线程ID与访问的内存地址应该保持单调递增关系。即thread 0访问地址0，thread 1访问地址1，以此类推。
我在项目中遇到过这样一个案例：一个矩阵转置操作，因为访问模式不对，性能只有理论峰值的5%。后来我把行访问改成了列访问，性能提升了15倍。嗯，你没看错，就是15倍。
// ❌ 非合并访问：每个线程访问不同行
float value = matrix[threadIdx.x * width + col]; // 跨步访问
// ✅ 合并访问：连续线程访问连续地址
float value = matrix[row * width + threadIdx.x]; // 连续访问警告：不要以为编译器会自动优化内存访问。我见过太多人写了非合并访问的代码，然后抱怨GPU性能不行。记住：合并访问是程序员的责任，不是编译器的责任。
3.4 知识体系总览
为了让你更直观地理解这些概念之间的关系，我画了一张图。这张图展示了我对GPU编程模型的理解框架。
这张图把三个核心概念放在了一起。你看，SIMT是执行模型，决定了指令怎么跑；线程层次结构是组织模型，决定了线程怎么分组；内存访问模式是优化模型，决定了数据怎么搬。三者缺一不可。
3.5 实战建议
说了这么多理论，最后给点实战建议吧。
- 先分析数据访问模式：写代码前，先想清楚数据是怎么被访问的。是连续访问还是随机访问？这决定了你的性能天花板。
- 合理配置Block大小：我一般用128或256个线程/Block。太小了硬件利用率低，太大了寄存器压力大。
- 避免warp发散：能用三元运算符就别用if-else。能用位运算就别用分支。
- 善用共享内存：共享内存比全局内存快两个数量级。但要注意bank conflict。
我的经验：刚开始学GPU编程时，我总想着把所有功能都塞进一个核函数。后来发现，拆分成多个小核函数，反而更容易优化。每个核函数只做一件事，做好一件事。
好了，这一章的内容就到这里。记住：GPU编程不是玄学，是有规律可循的。掌握了SIMT、线程层次结构和内存访问模式，你就掌握了GPU编程的三大支柱。
公众号:蓝海资料掘金营，微信deep3321