一、GPU架构基础：从CPU到GPU的思维转变

说实话，我第一次接触GPU编程时，脑子里全是CPU的那套思维。结果呢？写出来的代码跑得比CPU还慢。后来我才明白——GPU不是更快版本的CPU，它是另一种生物。

1.1 CPU vs GPU：设计哲学的根本差异

CPU的设计目标是低延迟。它要快速响应各种任务，所以有庞大的缓存、复杂的控制单元、强大的分支预测。说白了，CPU是个「全能选手」——什么活都能干，但一次只能干几件。

GPU的设计目标是高吞吐。它不在乎单个任务跑多快，它在乎的是同时处理成千上万个任务。GPU把大部分晶体管都给了计算单元，缓存和控制单元被压缩到极致。

我打个比方：CPU就像10个博士生，什么难题都能解；GPU就像10000个小学生，只会做加减法，但人多力量大。你让博士生去算1万次加法，他算得慢还嫌无聊；让小学生去算，瞬间搞定。

我在项目中遇到过一位同事，他花了两周优化一个串行算法，速度提升了3倍。我接手后用GPU重写，直接提速200倍。他问我秘诀，我说：「别跟CPU比单核速度，跟它比核数。」

1.2 SM（流多处理器）—— GPU的心脏

SM的全称是Streaming Multiprocessor，流多处理器。它是GPU最基本的计算单元。你可以把SM想象成一个小型计算中心，里面包含：

• CUDA核心：负责整数和浮点运算
• 共享内存：SM内部的高速缓存
• 寄存器文件：每个线程的私有存储
• 调度器：管理线程块的执行

一个GPU可能有几十到上百个SM。比如NVIDIA的A100有108个SM，每个SM包含64个CUDA核心。算一下：108 × 64 = 6912个核心。这就是为什么GPU能同时处理海量数据。

1.3 Warp—— GPU的调度单位

Warp是GPU执行指令的基本单位。一个Warp包含32个线程，它们执行完全相同的指令。这就是所谓的SIMT（单指令多线程）模型。

为什么会这样？因为GPU的设计哲学是「用控制单元的简化换取计算单元的增多」。32个线程共享一个指令单元，同时执行同一条指令，只是处理的数据不同。

if (threadIdx.x % 2 == 0) {
    // 偶数线程执行
    data[threadIdx.x] = a + b;
} else {
    // 奇数线程执行
    data[threadIdx.x] = a * b;
}

嗯，这里要注意：尽量让同一个Warp里的线程走相同的分支。如果实在避免不了，也要让分支粒度大于Warp大小（32）。

1.4 内存层次结构—— 速度与容量的博弈

GPU的内存层次和CPU类似，但更极端。从快到慢、从小到大排列：

你想想看，寄存器只要1个周期，全局内存要400个周期。这差距有多大？访问一次全局内存的时间，够寄存器干400次活了。

寄存器（Register）

每个线程私有的存储空间，速度最快。但数量有限——每个线程最多255个32位寄存器。如果你在核函数里声明太多局部变量，编译器会寄存器溢出，把数据挤到慢速的本地内存里。我见过一个案例，就因为多声明了几个double数组，性能掉了5倍。

共享内存（Shared Memory）

这是GPU编程的核心武器。同一个线程块内的所有线程可以共享数据，速度比全局内存快80倍。我常用的优化套路是：先把数据从全局内存搬到共享内存，让线程块内的线程协作处理，最后再写回全局内存。

__global__ void matrix_mul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    // 分块加载到共享内存
    As[ty][tx] = A[by * TILE_SIZE + ty][bx * TILE_SIZE + tx];
    Bs[ty][tx] = B[by * TILE_SIZE + ty][bx * TILE_SIZE + tx];
    __syncthreads();  // 同步，确保所有线程加载完成
    
    // 使用共享内存计算
    float sum = 0;
    for (int k = 0; k < TILE_SIZE; k++) {
        sum += As[ty][k] * Bs[k][tx];
    }
    C[by * TILE_SIZE + ty][bx * TILE_SIZE + tx] = sum;
}

注意那个__syncthreads()，它是个路障——确保所有线程都加载完数据后，才能开始计算。少了它，有的线程可能还在读数据，有的已经开始算了，结果全乱套。

本地内存（Local Memory）

这是个容易混淆的概念。名字叫「本地」，其实物理上在全局内存里。当寄存器不够用时，编译器会把变量溢出到本地内存。速度很慢，尽量避免。

怎么避免？少声明大数组、少用double、合理分配寄存器使用量。我习惯在编译时加--ptxas-options=-v参数，查看寄存器使用情况。如果发现寄存器溢出，就调整代码结构。

1.5 思维转变总结

从CPU到GPU，你需要改变三件事：

1. 从串行到并行：别想着怎么让一个线程跑得快，想怎么让一万个线程一起跑
2. 从缓存到共享内存：CPU靠缓存自动优化，GPU需要你手动管理数据搬移
3. 从分支到数据：CPU擅长复杂分支，GPU讨厌分支——尽量让数据决定计算路径

我记得第一次用GPU加速一个图像处理算法，没做任何优化就比CPU快了20倍。我当时很兴奋，觉得GPU编程不过如此。后来优化到200倍时，我才意识到——GPU的潜力远比你想象的大，但需要你理解它的思维方式。

这张图展示了GPU的内存层次。从底部的寄存器（最快、最小）到顶部的全局内存（最慢、最大）。编程的核心就是让数据尽量待在底层——多用寄存器、共享内存，少碰全局内存。

好了，这一章就到这里。记住：GPU编程不是写代码，是管理数据流。你控制好数据在哪里、怎么移动，性能自然就上来了。