第2章 GPU架构基础

第2章：GPU架构基础：流处理器（SM/CU）概念、线程束（Warp/Wavefront）、内存层次结构

好，咱们正式开始聊GPU的内部构造。说实话，很多初学者一上来就被各种缩写搞懵了——SM、CU、Warp、Wavefront……我当年也一样。记得我第一次看NVIDIA的架构白皮书，满脑子都是问号。但别急，这些东西说白了就是几个核心概念，捋顺了就好。

2.1 流处理器（SM/CU）——GPU的“计算小分队”

先说说SM（Streaming Multiprocessor），NVIDIA这么叫。AMD那边叫CU（Compute Unit）。名字不同，干的活一样——它们就是GPU里真正干活的核心单元。

你可以把SM想象成一个“计算小分队”。每个小分队里有很多“小兵”，也就是CUDA Core（NVIDIA）或者Stream Processor（AMD）。这些小兵负责执行具体的数学运算，比如加法、乘法、逻辑判断什么的。

我习惯用一个比喻：一个SM就像一个小型工厂车间。车间里有几十个工人（CUDA Core），有仓库（寄存器文件和共享内存），还有调度员（Warp Scheduler）。工人不能单独接活，必须组队干活。这个“队”就是后面要讲的线程束。

关键点：SM是硬件概念，线程是软件概念。一个SM可以同时管理成千上万个线程，但真正同时执行的，只有那么一小撮。

我在项目中遇到过一个问题：有人把线程块（Thread Block）设置得特别大，以为这样能压榨性能。结果呢？SM里的资源（寄存器、共享内存）被占满了，能同时运行的线程块数量反而变少了。嗯，这就是典型的“贪多嚼不烂”。

2.2 线程束（Warp/Wavefront）——GPU的“执行单位”

接下来是线程束。NVIDIA叫Warp，AMD叫Wavefront。名字不同，本质一样——它是GPU执行指令的最小单位。

为什么会这样？CPU可以一个线程一个线程地执行，GPU为什么非要搞个“束”？

说白了，GPU的设计哲学是“以量取胜”。它要同时处理成千上万个线程，如果每个线程都独立调度，调度器会疯掉。所以GPU把32个（NVIDIA）或64个（AMD）线程绑在一起，组成一个Warp。这32个线程执行同一条指令，只是处理不同的数据。这就是所谓的SIMT（单指令多线程）。

你想想看，这就像教官喊“向右转”，32个士兵同时向右转。指令是一样的，但每个士兵的初始朝向可能不同，所以转完后的结果也不同。这就是“同指令，不同数据”。

避坑指南：我曾经在写核函数时，让一个Warp里的线程走了不同的if-else分支。结果呢？性能直接腰斩。因为GPU会串行执行所有分支，最后再合并。这就是所谓的“线程束发散”（Warp Divergence）。记住：尽量让同一个Warp里的线程走相同的路径。

这里有个表格，帮你快速对比一下：

概念	NVIDIA	AMD	大小
执行单位	Warp	Wavefront	32 / 64 线程
计算核心	CUDA Core	Stream Processor	每个SM/CU多个
调度器	Warp Scheduler	Wavefront Scheduler	每个SM/CU多个

2.3 内存层次结构——GPU的“数据仓库”

内存这块，是GPU编程最容易出问题的地方。我见过太多人把数据放错地方，导致性能惨不忍睹。

GPU的内存是分层的，从快到慢、从近到远：

寄存器（Register）：最快，但每个线程只有那么一点点。我习惯说它是“贴身口袋”，放最常用的东西。
共享内存（Shared Memory）：同一个线程块内的线程可以共享。速度很快，但容量有限（通常几十KB）。我在项目中经常用它来做数据缓存，比如矩阵分块计算。
全局内存（Global Memory）：所有线程都能访问，但速度最慢。说白了就是GPU的“主存”，显存就是它。
常量内存（Constant Memory）：只读，有缓存，适合放那些不会变的数据，比如系数表。
纹理内存（Texture Memory）：专门为图形处理优化的只读内存，有空间局部性缓存。做图像处理时特别好用。

注意：全局内存的访问延迟大约是几百个时钟周期，而寄存器几乎是零延迟。所以，能用寄存器就别用共享内存，能用共享内存就别用全局内存。这个原则，我每次写核函数都会默念一遍。

这里有个简单的代码示例，展示怎么用共享内存来加速：

// 一个简单的矩阵乘法，使用共享内存做分块
__global__ void matMulShared(float *A, float *B, float *C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    int row = by * BLOCK_SIZE + ty;
    int col = bx * BLOCK_SIZE + tx;
    
    float sum = 0.0f;
    
    // 分块计算
    for (int k = 0; k < N / BLOCK_SIZE; ++k) {
        As[ty][tx] = A[row * N + (k * BLOCK_SIZE + tx)];
        Bs[ty][tx] = B[(k * BLOCK_SIZE + ty) * N + col];
        __syncthreads();  // 同步，确保数据加载完成
        
        for (int i = 0; i < BLOCK_SIZE; ++i) {
            sum += As[ty][i] * Bs[i][tx];
        }
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}

你看，这里用了__shared__关键字声明共享内存，然后用__syncthreads()做同步。为什么需要同步？因为同一个线程块里的线程需要等所有数据都加载到共享内存后，才能开始计算。不然读到的是脏数据。

我个人习惯在写这类代码时，先估算一下共享内存的用量。比如一个BLOCK_SIZE是16，那As和Bs各占16*16*4=1024字节，总共2KB。如果SM的共享内存是48KB，那可以同时跑24个这样的线程块？不对，还要考虑寄存器和线程数的限制。嗯，这里面的门道挺多的，后面章节会细讲。

2.4 小结

这一章我们聊了三个核心概念：

SM/CU：GPU的计算单元，里面有很多计算核心和调度器。
Warp/Wavefront：GPU执行指令的最小单位，32或64个线程绑在一起。
内存层次：从寄存器到全局内存，速度递减，容量递增。

记住一句话：GPU编程的本质，就是管理好你的线程束和内存访问模式。这两点做好了，性能不会差到哪里去。

下一章，我们会深入聊聊线程块和网格的配置，以及怎么根据硬件参数来优化。到时候我会分享一个我踩过的坑——因为线程块大小没选对，导致性能差了10倍。嗯，那是个挺惨痛的经历，但也是个很好的教训。

📚 课程目录