4、GPU架构与并行计算：SIMT模型、线程束调度、GPU内存层次
好，我们进入GPU这一块。说实话，CPU和GPU的思维方式完全不同。CPU像个精明的管家，什么杂活都能干，但一次只能处理几件事。GPU呢？它像个包工头，手下有几千个工人，干的是重复性的体力活——并行计算。
我个人习惯把GPU看作一个「吞吐量优先」的机器。它不在乎单个任务跑多快，它在乎的是：一秒钟能处理多少个任务。这个思路贯穿了整个GPU设计。
4.1 SIMT模型：单指令多线程
先聊SIMT。全称是Single Instruction, Multiple Threads。说白了就是：一条指令，驱动几十个线程同时干活。
你可能会问：这和SIMD有什么区别？嗯，区别大了。SIMD是数据并行，一条指令操作一组数据。SIMT是线程并行，一条指令控制一组线程，每个线程有自己的寄存器、程序计数器。线程之间可以走不同的分支——虽然代价很大。
核心要点：SIMT模型让程序员感觉自己在写多线程程序，但硬件层面其实是SIMD的变种。这种「假象」降低了编程难度，但代价是分支发散时的性能损失。
我在项目中遇到过一个问题：一个简单的if-else语句，让GPU性能直接腰斩。为什么？因为线程束（warp）里的32个线程必须执行相同的指令。如果一半走if，一半走else，那这两半就得串行执行。这就是所谓的「分支发散」。
避坑指南：我曾经在一个图像处理项目中，因为一个条件判断写得不小心，导致warp内线程频繁发散。后来我把条件判断改成了「掩码+选择」的方式，性能提升了40%。记住：尽量让同一个warp里的线程走相同的分支。
4.2 线程束调度：硬件如何管理千军万马
GPU里有成百上千个核心，但调度单元是「线程束」（warp）。NVIDIA的warp大小是32个线程，AMD的wavefront是64个线程。我个人习惯用NVIDIA的术语，毕竟用得最多。
调度器怎么工作？它维护一个warp池，每个warp有4种状态：
- 选中的（Selected）：正在执行指令
- 就绪的（Ready）：等待执行
- 停滞的（Stalled）：等待数据或资源
- 完成的（Completed）：执行完毕
调度器每周期从就绪的warp里选一个发射指令。如果某个warp因为访存延迟停滞了，调度器立刻切换到另一个就绪的warp。这就是「零开销上下文切换」——硬件级别的。
技巧：想要隐藏访存延迟？那就让每个SM（Streaming Multiprocessor）里驻留尽可能多的warp。我一般建议至少16个warp/SM，这样即使一半在等数据，另一半还能干活。
你想想看，CPU切换线程要保存恢复寄存器，几百个周期就没了。GPU呢？每个warp有自己的寄存器文件，切换就是改个指针，一个周期搞定。这就是为什么GPU能容忍高延迟——它用并行度来掩盖延迟。
4.3 GPU内存层次：从快到慢，从贵到贱
GPU的内存层次，说白了就是「谁离计算单元近，谁就快」。但和CPU不同的是，GPU的内存层次更强调带宽，而不是延迟。
我画了一张图，帮你理清这个层次结构：
这张图你看懂了吗？从上到下，速度越来越慢，容量越来越大。寄存器最快，但每个线程只有那么几十个。全局内存最慢，但容量最大。
4.3.1 寄存器
寄存器是GPU里最快的存储。每个线程私有，编译器会尽量把局部变量分配到寄存器。但寄存器数量有限——NVIDIA的每个SM有65536个32位寄存器，分给所有驻留线程。如果一个线程用太多寄存器，就会减少SM能同时容纳的线程数，降低并行度。
注意：我曾经调试一个kernel，发现占用率只有25%。查了半天，原来是一个临时数组被分配到了寄存器，每个线程用了128个寄存器。改成共享内存后，占用率升到75%，性能翻倍。寄存器不是越多越好，够用就行。
4.3.2 共享内存
共享内存是GPU编程的「杀手锏」。它位于芯片上，延迟只有几十个周期，带宽极高。同一个Block内的线程可以通过共享内存交换数据。
我建议你把它当作「手动管理的缓存」。为什么是手动？因为硬件不会自动帮你把数据搬进共享内存，你得自己写代码控制。但好处是——你知道数据在哪里，什么时候可用。
举个典型用法：矩阵分块乘法。把大矩阵切成小块，每个Block加载一块到共享内存，然后反复使用。这样全局内存的访问次数从O(N³)降到O(N²)，性能提升巨大。
4.3.3 全局内存
全局内存就是板载的DRAM，容量大但延迟高。访问全局内存时，有个关键概念叫「合并访问」（coalesced access）。什么意思？就是同一个warp里的32个线程，如果访问的地址是连续的，硬件会把32次访问合并成一次大的突发传输。效率极高。
反过来，如果每个线程访问的地址是随机分散的，那就得32次独立访问，带宽利用率惨不忍睹。
经验之谈：我优化过的一个项目，原本全局内存访问模式是「按列访问」，带宽利用率只有10%。改成「按行访问」后，利用率升到90%，性能提升了8倍。记住：让warp内的线程访问连续地址，这是GPU优化的第一课。
4.4 实战：一个简单的向量加法
说了这么多理论，来点实际的。下面是一个向量加法的CUDA代码，展示了线程组织和内存访问的基本模式：
// 向量加法 kernel
__global__ void vecAdd(float* A, float* B, float* C, int N) {
// 计算全局线程ID
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 边界检查
if (idx < N) {
C[idx] = A[idx] + B[idx];
}
}
// 调用方式
int N = 1 << 20; // 100万个元素
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);这段代码里，每个线程处理一个元素。threadIdx.x是线程在Block内的索引，blockIdx.x是Block的索引，blockDim.x是Block大小。三者组合得到全局唯一的线程ID。
注意那个边界检查。因为N不一定是Block大小的整数倍，最后一个Block可能有多余的线程。不加检查的话，这些线程会访问越界内存，导致程序崩溃或数据错误。
关键点：GPU编程的本质就是「把问题分解成大量细粒度任务，每个线程处理一小块」。线程之间尽量独立，减少同步和通信。如果必须通信，优先用共享内存。
嗯，GPU这块内容确实不少。SIMT模型让你用线程思维编程，线程束调度让硬件高效利用计算资源，内存层次则决定了你能跑多快。我个人觉得，理解这三者的关系，是掌握GPU编程的基石。
最后说一句：别被那些复杂的术语吓到。GPU说到底就是个「人多力量大」的架构。你只要记住——让线程尽量独立，让访存尽量连续，让资源尽量饱和。做到这三点，性能不会差。
我的习惯：每次写kernel之前，先画一张「数据流图」，标清楚数据从哪里来（全局内存）、在哪里处理（寄存器/共享内存）、最终写到哪里去。这样写出来的代码，结构清晰，性能也容易优化。
公众号:蓝海资料掘金营，微信deep3321