文章

GPU-Architecture

Architecture
#GPU
本文目录 13

GPU 架构

0.引言

GPU (Graphic Processing Unit)，原本作为图像处理的工具，为游戏的渲染服务，在AI时代因为其并行计算性而得到了大量的使用 Nvidia 搭建的在GPU上的编程平台 CUDA ，正是随着ai对于并行计算的需求增长而一飞冲天

最近看到 groq 使用的 LPU 带来的惊人的推理速度，想要了解一下它的架构，但是发现自己连现在的 GPU 架构也不是很了解，所以回炉重造一下

主要从两个方面入手：

  • GPU 的硬件架构
  • 编程模型，以及与硬件的对应
  • Memory设计

更新于26.5.6

1. 硬件架构（Hardware Architecture）

现在最广为使用，和编程结合的最紧密的还是有CUDA生态的Nvidia家的卡

而这些卡在不同代际之间的设计往往是不同的，更详细的可见演进历史

这里主要是基于现在流行的Hopper架构和Blackwell架构来写

单元分区

这是24年最新 blackwell 架构的硬件结构。3

从大到小，结构是这样的：

GPU -> GPC (Graphics Processing Cluster) 
    -> TPU (Texture Processing Cluster)
    -> SM  (Streaming Multiprocessor)

一张 GPU 最核心的单元就是 SM。可以看到下图中，一个 SM 中有 4 个子分区，每一个子分区都包含：

  • 控制单元
    • Warp Scheduler：从被分配的 Warp 中根据优先级和数据就绪状态调度 Warp
    • Dispatch Unit：把具体的指令分发给对应的计算单元
  • 计算单元
    • CUDA cores：执行最基础的标量运算
    • Tensor core：一个指令直接完成两个小矩阵的乘累加（MMA）
    • SFU (Special Funciton Unit)：处理一些非线性数学函数
    • LD/ST：负责从缓存/存储搬运数据
  • 缓存单元
    • L0 Instruction cache
  • 存储单元
    • Register

SM 示意图

hopper架构

hopper-SM

blackwell架构

5

在整个 SM 中，子单元相互之间会共用：

  • Share Memory
  • L1 cache
  • L1 Instruction Cache
  • Texture Unit

以上的这些有关 cache 的，在后面写

从硬件架构上，我们称之为SIMT架构，single instruction multiple threads

不难理解，就是一条指令会操纵多个线程，或者多个线程会运行同一条指令

TODO:TMA, Tensor Memory Accelerator

硬件化的数据异步加载机制

2. 编程模型（Programming Model）

CUDA core

CUDA core 是一个一个能进行标量运算的小核心

因此需要一种编程模型去让他们协同做一件事比如矩阵乘法

在编程模型中，由大到小的层级是

Grid -> Block -> Warp -> Thread
  • Thread：GPU最基本执行单元，每个线程执行同一个函数，但处理不同数据
  • Warp：
    • 1 Warp = 32 Thread
    • Warp 是调度的最小单位
    • 同一个 Warp 内必须执行相同指令
  • Block：
    • 由多个线程组成
    • 可以 1D / 2D / 3D
    • Block 内的线程：
      • 共享 Shared Memory / L1 cache
  • Grid：
    • 由多个 Block 组成
    • Grid 之间的不同的 Block:
      • 不能直接同步
      • 不共享 shared memory

image-20260929104242213

例如这张图表示了一个Grid，其中有2*3 = 6 个block；而blockdim.x=16, y=4;

而硬件在进行调度的时候，不是按照xy的坐标，而是按照一个一位的线性编号来调度的；比如这里一个[1,14]位置的编号就是1*16+14 = 30；而一个warp的分组就是按照这个线性编号来进行32倍数的划分的

既然都要转换为线性的，那为什么block要设计成一个二维的呢？

这样做代码的可读性会更好，也方便和矩阵对应起来；同时如果直接用一维的block去处理二维的矩阵的话，最后image[row][col] = ...;这一步难免要写row = tid / N，而整除和取模是更高级的运算指令，需要更多的周期来进行处理

硬件架构和编程模型的映射联系如下图correspond

一个 Grid 被扔给 GPU（或者说是一个GPC吧） 处理，其中的 Block 被分配到不同的 SM 处理，block 中的 warp 则是被分给 SM 的子单元，最终 Thread 由单个 CUDA core 处理

为了提升吞吐量，同一个SM往往会被分配到多个 Block，这些 Block 会被[分配一定的存储空间](# 3-memory-design)，他们内部的 Warp 都会接受调度并运行

GPU 遵循的是 SIMT（Single Instruction Multiple Threads模型，每个线程都有独立的寄存器，但是每个 Warp 内只执行相同的指令

举一个并行编程的例子：矩阵乘法

C=A×BC = A \times B

//CPU 的三层循环
for (int i = 0; i < M; i++)
  for (int j = 0; j < N; j++)
    for (int k = 0; k < K; k++)
      C[i][j] += A[i][k] * B[k][j];

并行粒度设计：

  • 一个线程负责计算 C 的一个元素
  • 一个 Block 负责计算一个子矩阵（tile）
  • 整个 Grid 覆盖整个输出矩阵

对于一个矩阵乘法来说，可以应用分块计算的方式，因此可以从结果矩阵[M,N]中分割出一些tile，对这些tile进行运算

假设 Block = 16 * 16 Thread，即一个 Block 分配有 8 个 Warp 。假设这个 Block 负责了一个 8 * 8的子矩阵（tile）；那么要让输出矩阵被完全覆盖，就需要 ceil(M/8) × ceil(N/8) 个 Block，这些Block组合在一起就是Grid

对于一个线程的索引，也就是它在矩阵中的真实的坐标，计算的顺序就是找到它在 Grid 里的第几个 Block，在这个 Block 里是第几个 Thread

row = blockIdx.y * blockDim.y + threadIdx.y
col = blockIdx.x * blockDim.x + threadIdx.x

得到索引后，Block 被分配给一个个 SM；每个 warp 执行同样的代码

float sum = 0.0f;

for (int k = 0; k < K; k++) {
    sum += A[row * K + k] * B[k * N + col];
}

C[row * N + col] = sum;

这样的设计称为是SPMD设计，Single-Program-Multi-Data

程序代码只有一个，但是它去操作不同的数据去完成了计算

Tensor Core

Tensor Core 是Warp级别的粒度，毕竟一个SM的子单元里只有一个tensor core

而每个时钟周期，该子单元的warp scheduler都只会选择一个warp来执行

因此他的编程抽象就是直接作为一个整体去调用

对于一个M=N=1024的结果矩阵，使用同样的block=256threads，CUDA core的逻辑就是这一个block里线程的数量就是对应的矩阵里元素的数量256；而block=256threads=8warps，一个warp调用tensorcore能负责一个16x16的tile，一个block能覆盖8x16x16=2048个元素

而warp对应的线程是根据线性id确定的，warp在block中的排列则是对应他要输出的矩阵的位置，以便于它写回结果

// 每个 Block 的 Warp 排列：4 行 × 2 列 = 8 个 Warp
#define WARPS_PER_BLOCK_M 4
#define WARPS_PER_BLOCK_N 2
#define WARPS_PER_BLOCK  (WARPS_PER_BLOCK_M * WARPS_PER_BLOCK_N)

// 每个 Block 覆盖的输出区域
#define BLOCK_TILE_M  (WARPS_PER_BLOCK_M * WMMA_M)  // 64
#define BLOCK_TILE_N  (WARPS_PER_BLOCK_N * WMMA_N)  // 32

// 每个 Block 的线程数：8 Warp × 32 线程 = 256
#define THREADS_PER_BLOCK (WARPS_PER_BLOCK * 32)

3. Memory Design

mem

这张图很好的显示了 GPU 的 memory architecture

Host (CPU DRAM)
        ↓ PCIe / NVLink
Global Memory (HBM / GDDR)
        ↓
L2 Cache (Chip-wide)
        ↓
SM 内部：
    ├─ L1 Cache
    ├─ Shared Memory
    ├─ Register File

SM 内部

Register (SRAM)

每个线程私有

当一个block被分派到一个sm执行的时候，根据编译器的规则会进行静态预留，从reg池子里划分出所有线程中专属的reg

之所以需要这么多物理寄存器，首先是因为gpu有海量的线程，根据硬件规则一个sm最多1024个线程，平均下来每个线程分配到16个寄存器；

其次是为了实现0开销切换。gpu不像cpu一样的，他为了最大化吞吐量，通过分配大量寄存器代替进行上下文的保存，从而节省了内存的读取成本。从空间资源上看，一个SM最多在前台同时运行4个warp（以这里只有四个子单元为例），其他没有运行的 warp 分配的reg是闲置的，但是这样允许了无开销的 warp reschedule，在一个warp需要等待上一步数据的时候，切换到可以运行的warp.而由于静态预留，这一步只需要修改指向reg的指针，而不需要重新malloc，从而提高吞吐效率

Shared Memory (SRAM)

  • 同一个 block 内所有的 Thread 共享
  • 由程序员手动管理
  • 可以并发访问

从硬件的角度来看，它的物理空间被等分为 32 个独立的Bank，而每个bank的行大小是4 byte；这块sRAM能容纳的行的大小就是 mem_size / 4 *32

shared memory 因为是手动分配的，可以理解为就是一个数组

L1 Cache (SRAM)

SM 内私有，自动缓存 global memory

与 shared memory 共享物理 SRAM

这玩意就和 CPU 上的缓存一样，是行缓存，一行的大小是128 byte；缓存的架构以及LRU的规则这里就不赘述了

总的下来这两个大概就是这样一种分配方式，中间的分割线就是硬件的一个截断位

SRAM

L2 Cache (SRAM)

全芯片共享，减少显存的访问压力

Global Memory

芯片外的 HBM / GDDR ，延迟相对来说高

Reference