如果你已经被“GPU程序跑得慢”、Nsight Profile 里一堆看不懂的指标或者“明明照着教程写的怎么还是快不起来”折磨过我猜你缺的可能不是一个更好的 API 写法而是一张关于 GPU 硬件本体的地图。网上讲 CUDA/ROCm 的教程一抓一大把但大多数只讲接口怎么调不讲芯片内部那套运行逻辑。恰恰是这套逻辑决定了你的代码能跑多快、显存会不会爆、为什么同样一段代码在别人的卡上飞快、在你的卡上却慢得离谱。这篇文章是 GPU 硬件原理架构系列的第一篇。我不打算上来就堆术语而是想顺一条线走清楚GPU 为什么被设计成如今这个样子、芯片里到底有哪些关键部件、几万个线程是怎么被组织起来“同时”执行的、数据在这些部件之间又是怎么流动的以及硬件事实如何反过来决定了 CUDA 和 ROCm 的标准写法。看完这篇你再去看那些“GPU 编程最佳实践”会发现很多规则可以自己推导出来而不是死记硬背。1. CPU和GPU的底层分歧延迟优先还是吞吐优先1.1 同样是芯片为什么设计思路差了这么远CPU 的设计目标非常明确让单条指令流执行得尽可能快。为了这个目标一颗现代 x86 处理器的面积里塞了大量“不直接算数”的电路——分支预测器、乱序执行引擎、重排序缓冲区、超大规模的多级缓存以及跨核心的一致性协议。这些电路存在的意义只有一个猜测程序下一步要干什么提前把数据准备好让那几个运算单元别闲着。单核里真正的 ALU/FPU 占比通常也就是 20% 到 30%。GPU 把目标换了不再追求单条指令多快而是追求单位时间里能执行多少条指令。所以 GPU 上那些“伺候单线程”的复杂控制逻辑全被砍掉换成大量简单的计算单元。以一块 NVIDIA 的 SM 为例里面有上百个 FP32 计算单元但几乎没有分支预测也没有乱序执行。整张卡积少成多可以同时跑几万甚至几十万条线程。行业里习惯用一句话概括CPU 是延迟导向GPU 是吞吐导向。CPU 擅长把单个任务尽快做完GPU 擅长把海量同类任务成批做完。如果你手头的问题是单任务线性逻辑GPU 不仅没有优势反而会因为更高的访存延迟和更低的单核频率跑得更慢。很多入门者误以为“GPU 核多就一定快”实际上只有任务具备足够的并行度GPU 的优势才会体现出来。1.2 晶体管预算、时钟和带宽全都在配合同一个目标任何芯片设计都绕不开“晶体管的钱花在哪”这个问题。一颗现代高端 CPU 的 L3 缓存动辄几十 MB加上分支预测器和乱序执行那套机构晶体管数量非常可观而一张旗舰 GPU比如 H100800 亿晶体管的绝大多数投给了计算单元和与之配套的数据通路。GPU 的 L2 缓存相对它的算力来说小得多这不是省成本而是设计哲学使然它默认你不会只靠一两个线程去跑而是有海量线程在同时运行。时钟频率上也能看出差异。消费级 CPU 睿频能干到 5GHz 以上GPU 核心频率通常只有 1.5GHz 到 2.2GHz。更低的主频换来了更高的流处理器数量和更可控的功耗因为 GPU 本来就是靠“宽度”而不是“高度”取胜。一块 GPU 上可能同时驻留几万个线程靠整体规模碾压单核的时钟优势。内存系统更是天差地别。CPU 的内存带宽依赖内存通道数量顶尖服务器 CPU 也就几百 GB/sGPU 走的是超宽显存总线HBM 这类堆叠显存能提供 3TB/s 以上的带宽。这一条对 GPU 太重要了——神经网络训练、渲染、科学计算全都是访存密集型任务GPU 能成为 AI 计算主力超高带宽功不可没。所以判断一张 GPU 值不值不能只看核心数和频率更要看它“喂得饱自己”的能力显存带宽、L2 带宽、寄存器文件总量以及片上互联。这也是后面几节要展开的内容。2. 宏观拆解GPU从GPC到SM芯片是分块管理的2.1 先看全局GPC/TPC/SM这套层级是怎么来的现代 NVIDIA GPU 的组织方式是分区块的几个 SM 组成一个 TPC几个 TPC 组成一个 GPC整颗芯片由多个 GPC 构成。以 Ampere 架构 GA102 为例7 个 GPC 分布在芯片四周中间是跨芯片共享的 L2 缓存和访存接口。AMD 这边结构类似Shader Engine 下挂若干 Compute UnitCU。叫法不同但分块管理的思想完全一致。分层不是软件工程师拍脑袋定的主要是三个原因。第一物理设计上芯片面积太大必须分成多个区域独立管理时钟和功耗否则局部发热和时钟偏差会失控。第二内存和缓存系统需要就近布局数据通路越短能效越好。第三调度需要合适的粒度把任务分配到 SM 之后SM 内部的资源管理才能做得精细。这个层级直接对应到 CUDA 的编程模型一个 kernel 启动后以 grid 形式存在grid 里的 thread block 会被调度到各个 SM 上一个 block 只能待在一个 SM 上不能拆分。所以 block 的大小和资源占用直接影响一个 SM 能同时容纳多少个 block也就影响整体占用率。调优时反复折腾 block 尺寸本质上就是在调整这张“任务分配表”。2.2 SM内部解剖一个车间里的关键工种把镜头拉近到 SM一个现代 SM 就像一个小车间里面有几种互相配合的“工种”。以 Ampere 架构为例一个 SM 包含 128 个 FP32 计算单元也就是大家常说的 CUDA core、若干 SFUSpecial Function Unit负责 sin/cos 等超越函数、若干 LSULoad/Store Unit负责访存指令还有 4 个 Tensor Core 和 4 个 Warp Scheduler。运算单元负责算术Tensor Core 专门做矩阵乘加从 Volta 开始引入是如今深度学习训练和推理的算力主力。Warp Scheduler 则负责不停地从就绪 warp 里选一条指令派发到执行单元SM 里有多个调度器可以并行派发。寄存器文件是 SM 里非常庞大的一块资源Ampere 一个 SM 有 65536 个 32 位寄存器合计 256KB比普通 CPU 的一个核心多得多。Shared Memory 和 L1 缓存共用一片片上 SRAM靠配置决定多少给缓存、多少给显式共享。这里的核心关系是SM 的并发能力由寄存器文件、Shared Memory、线程数上限共同决定三种资源任何一个先被占满都会限制驻留线程总量。这直接导致一个常见现象你以为限制性能的是“线程数”很多时候其实是寄存器或 Shared Memory 先满了。2.3 用手算推一遍为什么线程数和寄存器会打架这里给一个实际的计算例子。假设某个 SM 寄存器总量 65536 个你的 kernel 每个线程用 64 个寄存器那么一个 256 线程的 block 需要 256×6416384 个寄存器这个 SM 最多能驻留 65536/163844 个 block也就是 1024 个线程对应 SM 最大并发 2048 线程的 50% 占用率。如果每个线程用到 128 个寄存器一个 block 就要 32768 个寄存器SM 只能驻留 2 个 block也就是 512 线程占用率降到 25%。如果编译器发现寄存器不够会把部分局部变量溢出spill到本地内存——实际上就是显存——性能立刻崩。所以调整 block 大小、launch_bounds这些参数本质上就是在“线程并发数”和“每个线程可用寄存器数”之间找平衡。我的经验是不要迷信最大占用率先用 profiler 看 kernel 是访存瓶颈还是计算瓶颈再决定往哪个方向调。一个计算密集的 kernel占用率略低但每个线程手里寄存器充裕往往比强行拉高占用率导致溢出更快。3. Warp执行模型GPU性能最核心的“隐藏开关”3.1 32个线程绑成一个warp是指令派发的最小单元在 NVIDIA 硬件里32 个线程组成一个 warpSM 以 warp 为单位派发指令。一个 warp 里的 32 个线程执行同一条指令指令的操作数来自各自不同的寄存器这就是 SIMTSingle Instruction, Multiple Threads。AMD 里的叫法不同GCN/CDNA 时代叫 wavefront64 线程一组RDNA 后改成 32 线程的 wave32。名字不同原理一致。为什么是 32一个常见的解释是warp 太大分支分歧的代价会变高warp 太小取指与调度的开销占比不划算。同时 32 个线程读 32 个 4 字节的 float正好是 128 字节和缓存行粒度吻合非常适合做合并访存。这个数字从 Fermi 时代延续至今已经成为 CUDA 生态的基础假设大量优化技巧都是围着它转的。一个 256 线程的 block 会被切成 8 个 warp。warp 内的线程用 lane id 标识范围是 0 到 31。写 kernel 时经常看到 threadIdx 和 lane id 的换算关系但要注意 block 维度与 warp 切分方式不是简单的“行优先一眼看穿”真要吃透最好的办法是用 Nsight 或者写个小 kernel 打印 threadIdx 和 warp 的关系。3.2 延迟隐藏GPU不靠大缓存靠“人海战术”CPU 降低访存延迟靠缓存和乱序执行GPU 没有那么多缓存也不做乱序靠的是并发掩盖延迟。当某个 warp 执行的指令需要从显存取数它要等几百个周期才能拿到数据此时这个 warp 暂时派不上用场。不过 SM 里还有其他 warp调度器立刻切入一个可以执行的 warp让计算单元一直有事干。这个过程叫延迟隐藏latency hiding。要隐藏得彻底SM 里活动 warp 的数量得足够多。粗略估算一下如果显存访问延迟大约 600 个周期SM 每周期派发一条指令那么至少需要几百条“在飞行”的指令也就是要几十个 warp 同时在驻留才能把访存等待完全掩盖掉。这就是 occupancy占用率的意义活动 warp 数除以 SM 最大 warp 数。占用率太低延迟盖不住算力闲置占用率太高寄存器或 Shared Memory 不够反而触发溢出。实际调优时我通常先跑 Nsight看 Warp Stall 和 Occupancy 两个指标对比不同 block 大小下的表现而不是一上来就抄网上推荐的“256 线程”或“512 线程”。不同 kernel 的最优配置差别很大必须实测。3.3 分支分歧同一个warp同一时间只能走一条路因为 warp 是同步执行一条指令所以如果 warp 内线程出现分支分叉硬件就必须让它们依次执行每个分支。比如 if (tid % 2 0) 走 Aelse 走 B执行顺序就是先跑 A一半线程干活另一半被掩蔽也就是空转再跑 B另一半干活总耗时可能直接翻倍。这不是死锁也不是错误但它是一种不可见的浪费。对于性能敏感的 kernel尽量让 warp 内所有线程走同一条路径。比如处理循环边界条件时可以把边界处理的代码单独用一个 warp 或一个 block 处理而不是让每个 warp 都带一个 if。还有一个常用技巧很多条件判断可以用位运算、乘掩码等方式改写成算术操作从根源上避开分支分歧。理解这一点之后再看网上那些“用 min/max 代替 if”的写法就知道背后的原因了。4. 存储层级GPU的“血液循环系统”4.1 五级存储每一级负责什么GPU 的存储大致分五级访问速度从快到慢容量从大到小。理解这个结构是分析一切性能问题的起点。层级位置典型容量延迟量级说明寄存器文件SM 片上256KB/SM1~2 周期线程私有Shared Memory / L1SM 片上数十KB~百KB20~30 周期block 内共享可编程L2 缓存全芯片20~50MB200 周期左右跨 SM 共享显存 HBM/GDDR片外数十GB500~1000 周期主存主机内存 / PCIe系统系统内存微秒级数据交换寄存器文件在每个 SM 里都很夸张相当于给每个线程预分配了多个私人储物柜只要不用太多热点数据全都可以放里面。Shared Memory 和 L1 在物理上是同一块 SRAM靠配置决定多少给自动缓存、多少给显式共享。L2 是整颗芯片统一调度的最后一级片上缓存跨 SM 的数据共享和原子操作都要经过它。最后落到显存容量最大但速度掉了一个量级。4.2 内存合并GPU访存的第一法则GPU 访存的效率高度依赖 warp 内 32 个线程访问的地址是否连续。如果线程 tid 访问的地址是 base tid * 4每个线程读一个 4 字节的 float那么 32 个线程访问的是 128 字节连续区域硬件可以合并成一两次显存事务完成如果访问 base tid * 128每个线程间隔很大就会变成几十次独立事务有效带宽暴跌。这就是常说的内存合并memory coalescing它在任何 GPU 架构上都成立因为片外访存的代价实在太贵。优化的常见思路是调整数据结构布局把“按对象组织”改成“按字段组织”也就是 SoAStructure of Arrays而不是 AoSArray of Structures。举个例子一个粒子系统需要位置、速度、质量三个属性AoS 是每个粒子一个结构体SoA 是三个独立数组。GPU 线程要连续访问所有粒子的速度时SoA 天然合并AoS 则会出现跨步访问。这项改动在很多场景里能带来好几倍的访存加速而且改造成本通常不高。4.3 算力再高也怕带宽饿死Roofline视角GPU 算力增长一直快于显存带宽增长。拿 H100 举例FP32 算力大约 60 多 TFLOPSHBM3 显存带宽约 3.3TB/s。如果一条指令需要从显存读两个 FP32 数做一次加法理论刚需带宽等于算力乘以 8 字节远超显存能提供的量。所以绝大多数真实 kernel 都是访存受限的也就是常说的 memory-bound。判断一个 kernel 是 compute-bound 还是 memory-bound最通用的工具是 Roofline 模型纵轴是算力横轴是算术强度FLOP/Byte斜线部分是带宽限制区水平部分是算力限制区。先估算自己 kernel 的算术强度再和 GPU 的“FLOPS/带宽”比值对比就能知道优化方向。这个比值在 H100 上大约是 20 FLOP/Byte在消费级卡上通常更低。实际优化时如果程序跑在斜线区做再多的指令优化都没用得想办法减少访存量数据复用、缓存、合并访问、降低精度才是正路。5. 架构怎么倒逼出编程模型反推几条实战准则5.1 CUDA线程模型就是硬件模型的倒影CUDA 的 grid/block/thread 三层模型不是抽象出来的摆设它就是硬件执行层的直接映射一个 thread 对应一个 SIMT lane32 个 thread 组成一个 warp对应硬件派发单位一个 block 被调度到一个 SM 上内部再切成 warp整个 grid 横跨整张 GPU。理解这个对应关系很多“规定”就变得合理了。比如为什么 block 内可以用 shared memory 和 __syncthreads() 同步而 block 之间不行因为 shared memory 就在 SM 里block 里的线程都在同一个“车间”同步是片上操作跨 block 只能经过 L2 和全局内存代价高所以 CUDA 连原生 grid 级同步都很少见。再比如为什么性能建议里常说“相邻线程访问相邻地址”因为相邻线程就是同一个 warp 的相邻 lane它们的访存在硬件里天然可以被合并。这不是风格偏好是硬件事实。5.2 三条能直接上手的架构思维把前面的内容浓缩成三条实战判断遇到性能问题可以按这个顺序排查。第一先判断方向。一个 kernel 慢先看是卡在访存还是卡在计算。Nsight Compute 里看 SOLSpeed of Light相关的两个占比哪个接近 100% 就说明瓶颈在哪再决定优化动作。方向判断错了后面全是无用功。第二算资源账。每个线程用多少寄存器、多少 shared memory、block 开多大先用资源总量算一遍能驻留多少个 block、多少个 warp。很多时候性能差的根源就是寄存器 spill或者 shared memory 不够导致占用率太低。算完这笔账很多“玄学”调优立刻变成确定性工程。第三改数据布局优先于改指令。因为 GPU 普遍 memory-bound优化访存模式通常比优化算术指令收益大得多。把 AoS 改成 SoA加restrict帮助编译器做向量化用 float4 这类向量类型一次取 16 字节都是立竿见影的手段。这些写法的共同目的就是把“每次访存能搬回来的有用字节数”尽量拉高。我最初学 GPU 编程时也走了一段弯路拿着 CPU 的思维去优化在分支和循环上死磕结果收益甚微。后来把架构文档反复读了几遍对照 Nsight 面板一点点核对才意识到 GPU 调优的钥匙几乎全在硬件上。这个系列我计划按这个顺序写下去本篇是整体架构下一篇拆 warp 调度与指令流水线再往后是 Tensor Core 的矩阵乘原理、显存子系统与多卡互联。如果你也在折腾 GPU 相关的东西先把这一层的底子打牢后面讨论算子开发、深度学习训练的性能调优都会轻松很多。