GPU架构基础：SIMT执行模型、线程束调度、共享内存与全局内存

好，咱们今天聊聊GPU架构里最核心的几个概念。说实话，很多做AI加速器的朋友，一开始都容易把GPU当成“很多个CPU拼在一起”。嗯，这个理解其实不太对。GPU的思维方式，跟CPU完全不同。

SIMT执行模型：不是简单的SIMD

先说说SIMT，全称是Single Instruction, Multiple Threads。很多人一听，觉得这不就是SIMD吗？我刚开始也这么想。但实际用起来，差别大了去了。

SIMD是“一条指令，多个数据”，所有处理单元必须步调一致。而SIMT是“一条指令，多个线程”，每个线程有自己的程序计数器、寄存器和执行状态。说白了，每个线程看起来都像是一个独立的“小CPU”，但它们共享同一条指令流。

我在项目中遇到过这样一个坑：把CPU上的多线程代码直接搬到GPU上，以为能自动加速。结果发现，线程之间因为分支发散，性能反而下降了。为什么会这样？这就引出了下一个概念——线程束。

线程束调度：硬件如何“骗”你并行

GPU里最小的执行单位不是单个线程，而是线程束（Warp）。在NVIDIA的架构里，一个线程束通常是32个线程。这32个线程共享一条指令，但处理不同的数据。

你想想看，如果这32个线程都走同一个分支，那效率最高。但如果有人走if，有人走else，那就麻烦了。硬件会怎么做？它会先执行if分支，把走else的线程“屏蔽”掉，然后再反过来执行else分支。这就是所谓的“分支发散”。

线程束调度器是怎么工作的？它维护着一个就绪队列，里面放着可以执行的线程束。每个周期，调度器从中选出一个，发射一条指令。如果某个线程束在等待内存访问，调度器就切换到另一个线程束。这种“零开销切换”是GPU能隐藏内存延迟的关键。

共享内存与全局内存：谁快谁慢，心里要有数

GPU的内存层次，说白了就是“距离计算单元越近，越快，但越小”。咱们用一张表来对比一下：

内存类型	位置	大小	延迟	作用域
全局内存	片外DRAM	GB级别	几百个周期	所有线程
共享内存	片上SRAM	几十KB	几个周期	同一个线程块
寄存器	片上	每个线程有限	1个周期	单个线程

全局内存虽然大，但慢得让人着急。我习惯把全局内存比作“仓库”，东西多，但取一趟要跑半天。共享内存就像“工作台”，东西少，但随手就能拿到。

共享内存的妙用：数据复用与合并

共享内存最大的价值是什么？我个人觉得，是让同一个线程块里的线程可以“串通”数据。比如做矩阵乘法时，每个线程都需要读取同一行或同一列的数据。如果每个线程都去全局内存里取，那就重复读取了无数次。

正确的做法是：先把数据从全局内存搬到共享内存，然后所有线程从共享内存里读。这样全局内存的访问次数就大大减少了。

// 一个简单的共享内存使用示例
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float shared_A[TILE_SIZE][TILE_SIZE];
    __shared__ float shared_B[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    float sum = 0.0f;
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 协作加载数据到共享内存
        shared_A[ty][tx] = A[by * TILE_SIZE + ty][t * TILE_SIZE + tx];
        shared_B[ty][tx] = B[t * TILE_SIZE + ty][bx * TILE_SIZE + tx];
        __syncthreads();  // 同步，确保所有线程都加载完毕
        
        // 从共享内存计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += shared_A[ty][k] * shared_B[k][tx];
        }
        __syncthreads();  // 同步，防止覆盖数据
    }
    C[by * TILE_SIZE + ty][bx * TILE_SIZE + tx] = sum;
}

注意代码里的__syncthreads()。这个函数是线程块内的“集合点”，确保所有线程都执行到同一位置。我刚开始写GPU程序时，经常忘记加这个同步，结果数据还没加载完就开始计算，出来的结果全是错的。

全局内存的访问模式：合并访问

全局内存虽然慢，但如果访问模式对了，也能快不少。GPU的全局内存控制器有个特性：它会尽量把连续的线程访问合并成一次大的内存事务。

举个例子，如果线程0访问地址0，线程1访问地址4，线程2访问地址8……这样连续对齐的访问，硬件可以一次搞定。但如果线程0访问地址0，线程1访问地址100，那就得两次了。

我建议你在写代码时，始终让线程ID和内存地址保持线性关系。比如：

// 好的访问模式：连续
float val = data[threadIdx.x];

// 坏的访问模式：跨步
float val = data[threadIdx.x * 32];

嗯，这里要注意：跨步访问不仅浪费带宽，还会让缓存失效。我在做卷积优化时，就因为这个跨步问题，折腾了好几天才找到原因。

一张图看懂GPU内存与执行模型

下面我用一张SVG图来总结一下GPU的核心架构。这张图展示了线程、线程块、线程束、共享内存和全局内存之间的关系。

从这张图可以看得很清楚：全局内存是“大后方”，所有线程都能访问，但慢。共享内存是“前线指挥部”，同一个线程块内的线程可以快速交换数据。线程束是真正的执行单元，调度器负责在它们之间切换，隐藏延迟。

我个人觉得，理解这三者的关系，是写好GPU程序的关键。很多性能问题，归根结底就是没用好共享内存，或者全局内存访问模式不对。

好了，这一章的内容就到这里。记住这些基础概念，后面讲AI加速器的时候，你会发现很多设计思路都是从GPU这里借鉴过来的。