GPU 执行模型：SM、Warp、Block、Thread、SIMT浅析

概述

GPU 不是“一个大脑”，而是“很多个独立车间”：

GPU ├── SM 0（流式多处理器，一个独立车间） ├── SM 1 ├── ... └── SM 107（A100 共 108 个 SM） 每个 SM 内部： ├── 4 个 Warp 调度器 ├── 64 个 Warp 槽位（同时驻留 64 个 Warp） ├── 大量 CUDA Core / Tensor Core └── 192 KB L1 缓存 / 共享内存

执行时，程序员写的 Block 被分配到某个 SM

Block 内的 Thread 被硬件按 32 个一组打包成 Warp，Warp 以 SIMT 方式同步执行同一条指令。

Thread：最细粒度的工人

概述

每个线程执行同一段 Kernel 代码，但操作不同数据。

线程有自己的私有寄存器，用 threadIdx.x 区分“我处理哪个数据”。

数值例子：向量加法

A = [1, 2, 3, 4, 5, 6, 7, 8] B = [10, 20, 30, 40, 50, 60, 70, 80] C = A + B

启动 8 个线程，每个线程执行：

C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x];

线程	threadIdx.x	读取 A	读取 B	计算	写回 C
T0	0	1	10	1+10=11	C[0]=11
T1	1	2	20	2+20=22	C[1]=22
T2	2	3	30	3+30=33	C[2]=33
T3	3	4	40	4+40=44	C[3]=44
T4	4	5	50	5+50=55	C[4]=55
T5	5	6	60	6+60=66	C[5]=66
T6	6	7	70	7+70=77	C[6]=77
T7	7	8	80	8+80=88	C[7]=88

结果：

C = [11, 22, 33, 44, 55, 66, 77, 88]

关键： 8 个线程执行的是同一条指令C[i] = A[i] + B[i]，只是 i 不同。这就是 SIMT 的雏形。

Warp：32 个线程的“公交车”

概述

Warp 是 SM 调度的最小单元，固定 32 个连续编号的线程。Warp 内的 32 个线程必须同步执行同一条指令。

数值例子：Warp 内 32 线程同步执行

假设一个 Block 有 64 个线程，会被硬件分成 2 个 Warp：

Warp 0：线程 0~31 Warp 1：线程 32~63

执行C[i] = A[i] + B[i]时：

时刻 t1：Warp 0 的 32 个线程同时读 A[0~31]、B[0~31] 时刻 t2：Warp 0 的 32 个线程同时做加法 时刻 t3：Warp 0 的 32 个线程同时写 C[0~31] 时刻 t4：Warp 1 的 32 个线程同时读 A[32~63]、B[32~63]

32 个线程不是“一个接一个”，而是“一起动”。

这就是“公交车”比喻：32 个乘客同站同下。

Warp Divergence（线程束分化）

如果代码里有分支：

if (threadIdx.x < 16) { C[i] = A[i] + B[i]; } else { C[i] = A[i] - B[i]; }

Warp 0 的线程 0~15 走 if 分支，线程 16~31 走 else 分支。

但 Warp 内 32 个线程必须同步，硬件只能先执行 if 分支，再执行 else 分支，两个分支的线程互相等待。

数值例子：

线程 0~15：C[i] = A[i] + B[i] → 先执行 线程 16~31：C[i] = A[i] - B[i] → 后执行

原本 32 个线程并行，现在变成两批各 16 个串行，有效并行度减半。这就是 Warp Divergence 的代价。

避免方法： 尽量让同一个 Warp 内的线程走相同分支，比如按 threadIdx.x / 32 而不是 threadIdx.x % 2 来分支。

Block：一个施工队，独占一个 SM

概述

Block 是程序员指定的线程组，映射到一个 SM 上执行。

Block 内的线程可以共享共享内存、用__syncthreads()同步。

Block 内共享内存做归约求和

目标：对 1024 个数求和。

data = [1, 1, 1, ..., 1] （1024 个 1）

启动 1 个 Block，256 个线程，每个线程处理 4 个数：

__shared__ float sdata[256]; int tid = threadIdx.x; // 每个线程先把自己负责的 4 个数加起来 sdata[tid] = data[tid*4] + data[tid*4+1] + data[tid*4+2] + data[tid*4+3]; __syncthreads(); // 归约：256 -> 128 -> 64 -> ... -> 1 for (int s = 128; s > 0; s >>= 1) { if (tid < s) sdata[tid] += sdata[tid + s]; __syncthreads(); }

数值推演：
初始:

sdata[0~255] = 4, 4, 4, ..., 4 （每个线程负责 4 个 1，和为 4）

第 1 轮归约（s=128）：

sdata[0] = sdata[0] + sdata[128] = 4 + 4 = 8 sdata[1] = sdata[1] + sdata[129] = 8 ... sdata[0~127] = 8

第 2 轮（s=64）：

sdata[0~63] = 16

继续：

s=32: sdata[0~31] = 32 s=16: sdata[0~15] = 64 s=8: sdata[0~7] = 128 s=4: sdata[0~3] = 256 s=2: sdata[0~1] = 512 s=1: sdata[0] = 1024

最终 sdata[0] = 1024，与真实和一致。

关键：

• sdata 是 Block 内 256 个线程共享的，读写延迟远低于全局显存；
• __syncthreads() 保证所有线程写完 sdata 后，下一轮归约才读；
• 这个 Block 独占一个 SM 的共享内存和寄存器。

Block 不能跨 SM

一个 Block 的 256 个线程必须在同一个 SM 内执行，因为它们要共享 sdata。

不同 Block 之间不能直接共享共享内存，只能通过全局显存通信。

SM：独立车间

概述

SM 是 GPU 中独立调度和执行的基本硬件单元。A100 有 108 个 SM。

每个 SM 内部有：

• 4 个 Warp 调度器；
• 64 个 Warp 槽位（同时驻留 64 个 Warp）；
• 大量 CUDA Core / Tensor Core；
• 192 KB L1 缓存 / 共享内存。

数值例子：多 Block 分配到多 SM

假设：

网格：8 个 Block 每个 Block：256 个线程 A100：108 个 SM

硬件调度

Block 0 → SM 0 Block 1 → SM 1 Block 2 → SM 2 ... Block 7 → SM 7

每个 Block 的 256 个线程被分成：

256 / 32 = 8 个 Warp

所以：

SM 0 上驻留 Block 0 的 8 个 Warp SM 1 上驻留 Block 1 的 8 个 Warp

SM 同时驻留 64 个 Warp，意味着一个 SM 可以同时容纳多个 Block。

例如：

SM 0 上可以同时驻留： Block 0 的 8 个 Warp Block 8 的 8 个 Warp Block 16 的 8 个 Warp ... 总共不超过 64 个 Warp

Warp 调度器在 64 个 Warp 之间快速切换，某个 Warp 等内存时，另一个 Warp 立刻上计算单元，隐藏延迟。

数值例子：Warp 调度隐藏延迟

设：

• 一个 Warp 执行一次加法需要 4 个周期；
• 访问全局显存需要 400 个周期。

如果只有 1 个 Warp：

4 周期计算 → 400 周期等内存 → 4 周期计算 → ... 利用率 = 4 / (4+400) ≈ 1%

如果有 64 个 Warp：

Warp 0 等内存时，Warp 1 上计算 Warp 1 等内存时，Warp 2 上计算 ... 64 个 Warp 轮流上，计算单元几乎不空闲 利用率接近 100%

这就是 SM 同时驻留大量 Warp 的意义：用并行隐藏延迟。

SIMT：单指令多线程

概述

SIMT 是 NVIDIA GPU 的执行模型：一条指令广播给 Warp 内 32 个线程，每个线程操作自己的数据。

数值例子：SIMT 执行C[i] = A[i] * 2 + B[i]

Warp 0，32 个线程，i = threadIdx.x。

指令 1：加载 A[i]

32 个线程同时读 A[0]~A[31]

指令 2：乘以 2

32 个线程同时算 A[0]*2, A[1]*2, ..., A[31]*2

指令 3：加载 B[i]

32 个线程同时读 B[0]~B[31]

指令 4：相加

32 个线程同时算 A[i]*2 + B[i]

指令 5：写回 C[i]

32 个线程同时写 C[0]~C[31]

数值：
设:

A[0~31] = [1, 2, 3, ..., 32] B[0~31] = [10, 20, 30, ..., 320]

则:

C[0] = 1*2 + 10 = 12 C[1] = 2*2 + 20 = 24 C[2] = 3*2 + 30 = 36 ... C[31] = 32*2 + 320 = 384

32 个线程执行了5 条指令，完成了 32 个元素的 3 次运算。

如果串行，需要 32×3=9632×3=96 次运算；

SIMT 用 5 条指令覆盖，指令级并行度 32 倍。

SIMT vs SIMD

对比项	SIMD（CPU 向量指令）	SIMT（GPU）
指令	一条指令操作固定宽度向量	一条指令广播给 32 个线程
数据	寄存器中的向量	每个线程私有寄存器
分支	不支持或需特殊处理	支持，但有 Warp Divergence
编程	需显式向量化	按线程写，硬件自动打包

SIMT 的好处是程序员按“每个线程处理一个元素”来写，硬件自动把 32 个线程打包成 Warp，屏蔽了向量化细节。

一个完整 Kernel 的执行

Kernel

__global__ void vec_add(float* A, float* B, float* C, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { C[i] = A[i] + B[i]; } }

启动配置

N = 1000 blockDim.x = 256 （每个 Block 256 个线程） gridDim.x = ceil(1000/256) = 4 （4 个 Block）

线程分配

Block 0：线程 0~255 → 处理 i = 0~255 Block 1：线程 0~255 → 处理 i = 256~511 Block 2：线程 0~255 → 处理 i = 512~767 Block 3：线程 0~255 → 处理 i = 768~1023（但 N=1000，i=1000~1023 越界）

Warp 划分

每个 Block 256 线程 = 8 个 Warp：

Block 0：Warp 0（线程0~31）、Warp 1（32~63）、...、Warp 7（224~255）

越界处理

Block 3 的线程处理 i = 768~1023，但 N=1000，所以：

i = 768~999：正常执行 i = 1000~1023：if (i < N) 为假，跳过

Block 3 的 Warp 7（线程 224~255，对应 i=992~1023）：

线程 224~231（i=992~999）：执行 C[i] = A[i] + B[i] 线程 232~255（i=1000~1023）：不执行

这个 Warp 内出现了分支，但只是“部分线程活跃”，不是两个分支都执行，代价较小。

执行时间估算

设：

• 每个 Warp 执行一次加法需要 4 周期；
• 总 Warp 数：4 Block × 8 Warp = 32 个 Warp；
• A100 有 108 个 SM，4 个 Blocks 可以分到 4 个 SM 并行。

如果 4 个 Block 分到 4 个 SM：

每个 SM 处理 8 个 Warp 8 个 Warp 可以并行驻留，轮流上计算单元 总时间 ≈ 8 × 4 = 32 周期（粗略）

如果只有 1 个 SM：

32 个 Warp 驻留，轮流上计算单元 总时间 ≈ 32 × 4 = 128 周期

结论： Block 越多、SM 越多，并行度越高。A100 的 108 个 SM 能同时处理大量 Block，这是 GPU 海量并行的硬件基础。

总结

概念	是什么	作用	数值例子
Thread	最细粒度执行单元	每个线程处理一个数据元素	8 个线程做向量加法，c[i]=a[i]+b[i]
Warp	32 个连续线程的固定小组	SM 调度的最小单元，同步执行同一条指令	32 个线程同时读 A[0~31]，分支时分化
Block	程序员指定的线程组	映射到一个 SM，共享共享内存，可同步	256 线程用 sdata 归约求和，1024 个 1 求和得 1024
SM	独立调度和执行的硬件单元	驻留 64 个 Warp，4 个调度器，隐藏延迟	8 个 Block 分到 8 个 SM，每 SM 8 个 Warp
SIMT	单指令多线程执行模型	一条指令广播给 Warp 内 32 线程，各操作各的数据	5 条指令完成 32 个元素的 3 次运算

一句话： GPU 的执行模型是“SM 管 Block，Block 含 Thread，Thread 打包成 Warp，Warp 以 SIMT 方式同步执行”。

Thread 是最细的工人，Warp 是 32 人一组的公交车，Block 是独占一个车间的施工队，SM 是独立车间，SIMT 是“一声令下，32 人同时干不同活”的执行方式。

A100 的 108 个 SM、每 SM 64 个 Warp，共同构成了海量并行的硬件基础。