第二章：GPU架构基础

各位同学，今天我们来聊聊GPU最核心的架构概念。说实话，我刚入行时也被这些术语搞得晕头转向——SIMT、Warp、SM、各种内存……但搞懂了这些，你就掌握了GPU设计的钥匙。

2.1 SIMT执行模型：GPU的灵魂

SIMT，全称是Single Instruction Multiple Threads。说白了，就是一条指令控制多个线程同时干活。

你可能会问：这和CPU的SIMD有什么区别？嗯，区别大了。CPU的SIMD是程序员手动把数据打包，比如用AVX指令一次处理8个float。但GPU的SIMT是硬件自动管理的——你写一个线程的行为，硬件帮你复制成几十上百个线程一起跑。

关键点：SIMT模型下，每个线程有自己的程序计数器（PC）和寄存器状态。虽然它们执行同一条指令，但可以走不同的分支路径。当然，代价就是分支发散时性能会下降。

我在项目中遇到过这样一个坑：写了一个if-else分支，结果一半线程走if，一半走else。硬件不得不串行执行两个分支，性能直接腰斩。所以，写GPU代码时尽量让同一个Warp里的线程走相同路径。

2.2 Warp / Wavefront：硬件调度的基本单位

Warp是NVIDIA的叫法，AMD叫Wavefront。本质上是一回事——一组同时执行的线程。

NVIDIA的Warp大小是32个线程，AMD的Wavefront是64个线程。为什么是32或64？这跟硬件设计有关。我个人习惯把Warp看作一个“原子执行单元”——调度器每次调度一个Warp，而不是单个线程。

来看一个简单的例子：

// 假设我们有128个线程，Warp大小=32
// 硬件会把这128个线程分成4个Warp
// Warp 0: 线程0-31
// Warp 1: 线程32-63
// Warp 2: 线程64-95
// Warp 3: 线程96-127

__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];  // 同一个Warp内的线程执行相同指令
    }
}

避坑指南：我曾经在设计调度器时，忽略了Warp的上下文切换开销。每个Warp需要保存32份寄存器状态，如果寄存器压力大，切换成本会很高。后来我学会了在编译器层面做寄存器分配优化。

2.3 流多处理器（SM / CU）：GPU的计算核心

SM（Streaming Multiprocessor）是NVIDIA的叫法，AMD叫CU（Compute Unit）。它是GPU里真正干活的计算单元。

一个SM包含：

  • 多个CUDA核心（也叫SP，Streaming Processor）——执行算术运算
  • 共享内存（Shared Memory）——低延迟、高带宽的片上存储
  • 寄存器文件（Register File）——每个线程私有的最快存储
  • Warp调度器——负责选择和发射Warp
  • 加载/存储单元（LD/ST Unit）——处理内存访问
  • 特殊函数单元（SFU）——处理sin、cos、sqrt等

我记得有一次做性能分析，发现某个kernel的occupancy（占用率）很低。查了半天，原来是每个线程用了太多寄存器，导致SM能同时容纳的Warp数量减少。说白了，寄存器是有限资源，你用多了别人就进不来。

组件 功能 设计注意事项
CUDA核心 执行整数/浮点运算 注意流水线深度和发射宽度
Warp调度器 每个周期选择就绪的Warp发射 需要足够的就绪Warp来隐藏延迟
共享内存 同一Block内线程共享 Bank冲突会影响带宽
寄存器文件 每个线程私有 寄存器溢出会掉到Local Memory

2.4 内存层次结构：速度与容量的博弈

GPU的内存层次，说白了就是一套“速度越快、容量越小、价格越贵”的体系。从上到下依次是：

2.4.1 寄存器（Register）

最快，没有之一。每个线程独享，延迟1个周期。但数量有限——NVIDIA的SM通常有65536个32位寄存器，分给所有线程用。你想想看，如果每个线程用32个寄存器，那一个SM最多只能同时容纳2048个线程。

注意：我曾经见过一个kernel，因为寄存器压力太大，编译器把部分变量“溢出”到了Local Memory。性能直接从100%掉到30%。所以，写代码时尽量少用局部变量，或者用__launch_bounds__限制每个线程的寄存器数。

2.4.2 共享内存（Shared Memory）

位于SM内部，延迟大约20-30个周期。同一Block内的线程可以互相访问。典型大小是48KB或96KB（取决于架构）。

共享内存最大的坑是Bank冲突。共享内存被分成32个Bank（对应Warp大小），如果多个线程同时访问同一个Bank的不同地址，就会发生冲突，导致串行化。

// 好的访问模式：无Bank冲突
__shared__ float data[32][32];
float val = data[threadIdx.x][threadIdx.y];  // 每个线程访问不同行

// 坏的访问模式：16路Bank冲突
float val = data[threadIdx.y][threadIdx.x];  // 同一列，不同行

2.4.3 全局内存（Global Memory）

最大的存储空间，但也是最慢的。延迟400-800个周期。所有线程都能访问，但必须通过缓存。

访问全局内存时，要尽量满足“合并访问”（Coalesced Access）——同一个Warp的线程访问连续地址。否则，一次内存事务只能传输少量有用数据，浪费带宽。

2.4.4 本地内存（Local Memory）

这是个容易混淆的概念。名字叫“本地”，实际上存在全局内存里。当寄存器不够用时，编译器会把变量 spill 到这里。延迟和全局内存一样高。

如何判断变量被 spill 了？看PTX代码或者用profiler。我建议你在设计RTL时，给寄存器文件留足够的端口，减少spill的概率。

2.5 总结与个人体会

讲到这里，我想分享一点个人经验。做GPU架构设计，最难的不是理解这些概念本身，而是理解它们之间的相互影响。比如：

  • Warp大小影响调度器设计
  • 寄存器数量影响occupancy
  • 共享内存大小影响Block的并行度
  • 全局内存带宽影响整体吞吐

你想想看，这些参数环环相扣。改一个，其他都得跟着调。我当年设计第一代SM时，就是因为没算清楚寄存器文件和Warp调度器的关系，导致芯片面积超标30%。后来重新做floorplan才搞定。

下一章我们会深入Warp调度器的RTL实现，到时候你会看到这些概念是如何落地成具体电路的。

← 上一章 📖 返回目录 下一章 →