概述GPU 不是“一个大脑”而是“很多个独立车间”GPU ├── SM 0流式多处理器一个独立车间 ├── SM 1 ├── ... └── SM 107A100 共 108 个 SM 每个 SM 内部 ├── 4 个 Warp 调度器 ├── 64 个 Warp 槽位同时驻留 64 个 Warp ├── 大量 CUDA Core / Tensor Core └── 192 KB L1 缓存 / 共享内存执行时程序员写的 Block 被分配到某个 SMBlock 内的 Thread 被硬件按 32 个一组打包成 WarpWarp 以 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计算写回 CT0011011011C[0]11T1122022022C[1]22T2233033033C[2]33T3344044044C[3]44T4455055055C[4]55T5566066066C[5]66T6677077077C[6]77T7788088088C[7]88结果C [11, 22, 33, 44, 55, 66, 77, 88]关键 8 个线程执行的是同一条指令C[i] A[i] B[i]只是 i 不同。这就是 SIMT 的雏形。Warp32 个线程的“公交车”概述Warp 是 SM 调度的最小单元固定 32 个连续编号的线程。Warp 内的 32 个线程必须同步执行同一条指令。数值例子Warp 内 32 线程同步执行假设一个 Block 有 64 个线程会被硬件分成 2 个 WarpWarp 0线程 0~31 Warp 1线程 32~63执行C[i] A[i] B[i]时时刻 t1Warp 0 的 32 个线程同时读 A[0~31]、B[0~31] 时刻 t2Warp 0 的 32 个线程同时做加法 时刻 t3Warp 0 的 32 个线程同时写 C[0~31] 时刻 t4Warp 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~15C[i] A[i] B[i] → 先执行 线程 16~31C[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 个 Block256 个线程每个线程处理 4 个数__shared__ float sdata[256]; int tid threadIdx.x; // 每个线程先把自己负责的 4 个数加起来 sdata[tid] data[tid*4] data[tid*41] data[tid*42] data[tid*43]; __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 轮归约s128sdata[0] sdata[0] sdata[128] 4 4 8 sdata[1] sdata[1] sdata[129] 8 ... sdata[0~127] 8第 2 轮s64sdata[0~63] 16继续s32: sdata[0~31] 32 s16: sdata[0~15] 64 s8: sdata[0~7] 128 s4: sdata[0~3] 256 s2: sdata[0~1] 512 s1: 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 Core192 KB L1 缓存 / 共享内存。数值例子多 Block 分配到多 SM假设网格8 个 Block 每个 Block256 个线程 A100108 个 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 个 WarpSM 同时驻留 64 个 Warp意味着一个 SM 可以同时容纳多个 Block。例如SM 0 上可以同时驻留 Block 0 的 8 个 Warp Block 8 的 8 个 Warp Block 16 的 8 个 Warp ... 总共不超过 64 个 WarpWarp 调度器在 64 个 Warp 之间快速切换某个 Warp 等内存时另一个 Warp 立刻上计算单元隐藏延迟。数值例子Warp 调度隐藏延迟设一个 Warp 执行一次加法需要 4 个周期访问全局显存需要 400 个周期。如果只有 1 个 Warp4 周期计算 → 400 周期等内存 → 4 周期计算 → ... 利用率 4 / (4400) ≈ 1%如果有 64 个 WarpWarp 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 032 个线程i threadIdx.x。指令 1加载 A[i]32 个线程同时读 A[0]~A[31]指令 2乘以 232 个线程同时算 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 38432 个线程执行了5 条指令完成了 32 个元素的 3 次运算。如果串行需要 32×39632×396 次运算SIMT 用 5 条指令覆盖指令级并行度 32 倍。SIMT vs SIMD对比项SIMDCPU 向量指令SIMTGPU指令一条指令操作固定宽度向量一条指令广播给 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但 N1000i1000~1023 越界Warp 划分每个 Block 256 线程 8 个 WarpBlock 0Warp 0线程0~31、Warp 132~63、...、Warp 7224~255越界处理Block 3 的线程处理 i 768~1023但 N1000所以i 768~999正常执行 i 1000~1023if (i N) 为假跳过Block 3 的 Warp 7线程 224~255对应 i992~1023线程 224~231i992~999执行 C[i] A[i] B[i] 线程 232~255i1000~1023不执行这个 Warp 内出现了分支但只是“部分线程活跃”不是两个分支都执行代价较小。执行时间估算设每个 Warp 执行一次加法需要 4 周期总 Warp 数4 Block × 8 Warp 32 个 WarpA100 有 108 个 SM4 个 Blocks 可以分到 4 个 SM 并行。如果 4 个 Block 分到 4 个 SM每个 SM 处理 8 个 Warp 8 个 Warp 可以并行驻留轮流上计算单元 总时间 ≈ 8 × 4 32 周期粗略如果只有 1 个 SM32 个 Warp 驻留轮流上计算单元 总时间 ≈ 32 × 4 128 周期结论 Block 越多、SM 越多并行度越高。A100 的 108 个 SM 能同时处理大量 Block这是 GPU 海量并行的硬件基础。总结概念是什么作用数值例子Thread最细粒度执行单元每个线程处理一个数据元素8 个线程做向量加法c[i]a[i]b[i]Warp32 个连续线程的固定小组SM 调度的最小单元同步执行同一条指令32 个线程同时读 A[0~31]分支时分化Block程序员指定的线程组映射到一个 SM共享共享内存可同步256 线程用 sdata 归约求和1024 个 1 求和得 1024SM独立调度和执行的硬件单元驻留 64 个 Warp4 个调度器隐藏延迟8 个 Block 分到 8 个 SM每 SM 8 个 WarpSIMT单指令多线程执行模型一条指令广播给 Warp 内 32 线程各操作各的数据5 条指令完成 32 个元素的 3 次运算一句话 GPU 的执行模型是“SM 管 BlockBlock 含 ThreadThread 打包成 WarpWarp 以 SIMT 方式同步执行”。Thread 是最细的工人Warp 是 32 人一组的公交车Block 是独占一个车间的施工队SM 是独立车间SIMT 是“一声令下32 人同时干不同活”的执行方式。A100 的 108 个 SM、每 SM 64 个 Warp共同构成了海量并行的硬件基础。