最近这阵子团队准备自建训练集群方案评审时经常被问到同一个问题为什么花这么多钱买GPU反而训练吞吐上不去我意识到很多做AI Infra的同事包括一些已经写了好几年训练脚本的人对GPU架构的理解还停留在很多个小核心并行算这个层面。这是个挺要命的事。AI Infra这个岗位本质上是在跟硬件打交道你不懂硬件底层的运作逻辑就没法解释性能瓶颈到底出在哪也没法判断该砸钱买卡还是调代码。上一篇文章聊了AI Infra的整体框架这篇我专门把GPU架构这块掰开了讲清楚。别怕我不跟你扯一堆微架构代码名我会从最底层的设计哲学讲起把SM、显存带宽、Tensor Core这些词背后的门道全部说透。1. CPU和GPU的本质分歧一个为快而生一个为多而生聊GPU架构之前必须先搞清楚一个底层问题为什么我们不能直接用CPU来做AI训练现在CPU的核心数量也在涨动不动几十个核频率也高跑个推理任务也没见慢到哪去。答案在于CPU和GPU的设计目标从根本上就是相反的。1.1 CPU设计哲学低延迟CPU是为延迟敏感的任务设计的。操作系统、数据库、事务处理这类任务的特点是逻辑复杂、分支多、依赖性强。你执行完一条指令下一条指令的结果可能依赖于上一条的计算结果。这时候处理器最需要的是把单条指令跑得飞快把等待时间压到极限。为了做到低延迟CPU动用了大量晶体管做乱序执行、分支预测、大容量缓存。你打开任务管理器看到CPU占用率飙到100%觉得它在疯狂干活实际上里面有大量晶体管正在做流水线调度和分支猜测真正算数的单元占比并不高。一句话总结CPU的取向它更在乎单个人跑得快不快而不是同时能塞进去多少人。1.2 GPU设计哲学高吞吐GPU面对的从来不是这种任务。图形渲染里成千上万个顶点和像素需要被独立地做坐标变换、着色、光照计算。每个计算点之间互不依赖完全可以并行处理。AI训练的核心运算——矩阵乘法——也是同理一个32x32的矩阵乘法每个输出元素的计算互相独立不需要谁等谁。GPU把这些不需要互相依赖的计算任务堆给成千上万个简单核心去跑。它不追求单核速度追求的是单位时间内完成的计算总量也就是吞吐量。你可以把CPU想象成一个解题速度极快、但一次只能接一道题的专家把GPU想象成一千个中学毕业生同时开工每个人速度一般但一千道题一起做整体效率远远甩开专家。这也是为什么AI训练几乎被GPU垄断——矩阵乘法和卷积运算天然是高吞吐友好型任务。1.3 从Flynn分类法看两者定位计算机体系结构里有个经典的Flynn分类按指令流和数据流把计算系统分成四类CPU是典型的MISD多指令单数据的灵活变体GPU则是SIMTSingle Instruction, Multiple Threads单指令多线程的代表。SIMT的含义是同一时刻所有核心执行同一条指令但作用于不同的数据上。放到训练场景里就是一个warp线程束典型32个线程里的32个线程执行同样的乘加操作只不过每个线程处理矩阵中不同的元素。这种设计让GPU不需要像CPU那样做复杂的分支预测和乱序调度把晶体管预算全部砸在计算单元上。注意SIMT和CPU的SIMD不是一回事。SIMD是高层次的一条指令直接操作一组数据编程模型是显式向量化的。SIMT则是硬件层面把多个线程打包成组程序员视角下依然是普通线程模型实际执行时才被捆绑到同一调度器上。这就是GPU架构的总纲。下面所有的细节SM、显存、Tensor Core、NVLink都是这个基调下的展开。2. 核心组成拆解SM、Warp与执行管线理解了为吞吐而生之后我们进入GPU内部看看它是怎么把这套吞吐哲学落地的。2.1 SMGPU里的计算集群NVIDIA GPU里最核心的计算单元叫SMStreaming Multiprocessor流式多处理器。你可以把SM理解成一个小型的计算集群里面有自己的执行单元、寄存器文件、共享内存和调度器。一个GPU里通常有几十到上百个SM比如A100有108个SMH100有132个SMH800同样132个SM。每个SM内部又包含若干CUDA Core确切地说是FP32单元不过是通用叫法、Tensor Core张量核心后文细讲、特殊函数单元SFU以及负责调度这些单元的调度器。看一台GPU的算力表面上是看FP32 TFLOPS这个数字本质上是看SM的数量乘以每个SM每周期能执行的运算量。A100的FP32算力19.5 TFLOPSH100是67 TFLOPS带稀疏性数字差距背后就是SM规模和每周期执行能力的双重差异。2.2 WarpGPU执行调度的基本单位这是理解GPU架构的一个关键环节。GPU虽然号称几千个核心但它不是让每个核心独立执行各自指令的而是把32个线程绑成一个warp由调度器统一派发指令。为什么是32这是寄存器文件带宽和指令获取逻辑共同决定的硬件设计方案。32个线程组成的一个warp执行同一指令、处理不同数据这就是前面说的SIMT。warp的代价是什么如果32个线程里只有1个线程活跃其他31个都在等数据或者分支没走到那这个warp的执行效率就只剩1/32。这种浪费叫warp divergence线程束分化。矩阵乘法这类形态规整的运算几乎没有divergence所以GPU跑得飞起但如果你写代码时有大量if-else分支比如if (index % 2 0) {...} else {...}性能会断崖式下降。所以GPU编程有个原则不要让线程产生严重分化让所有线程做一致性的事情。这也是为什么很多手写kernel性能跑不过cuBLAS——人家已经把数据排布和线程形态优化到极致了。2.3 寄存器与Local Memory的陷阱每个SM有数量有限的寄存器典型是64K个32位寄存器。当你启动一个kernel时CUDA会按每个线程需要多少寄存器来分配。比如你一个kernel里声明了较多变量每个线程需要64个寄存器那么一个SM上最多同时驻留64K/641024个线程。如果寄存器不够用会发生什么变量会溢出spill到Local Memory——逻辑上是线程私有的物理上却躺在显存里。这就是性能杀手。一旦发生寄存器溢出看起来你的内存访问在本地实际上每次操作都要走显存带宽而显存带宽虽然大但延迟远比寄存器访问高。很多调优案例里一个kernel仅因寄存器溢出性能就掉了好几倍。你在Nsight Compute里看到local load/store计数不为0基本就是这个问题。所以调优的第一课永远是——检查寄存器使用量用__launch_bounds__限制每线程寄存器数量或者拆分kernel让每个线程的任务更轻盈。3. Tensor Core为什么这东西改变了AI计算格局如果说CUDA Core是常规计算单元那Tensor Core就是为AI量身定制的加速器。它的出现把AI算力推进了一个数量级。3.1 矩阵乘法的本质是乘累加AI计算的核心运算是矩阵乘法拆到微架构层面就是大量c a * b c的乘累加操作FMAfused multiply-add。一个32x32x32的矩阵乘法需要32768次乘累加。在老的GPU架构上比如P100、V100之前的通用CUDA Core这些乘法需要拆成普通指令逐条执行。即使有FMA指令每个核心每周期也只能做一次32位浮点乘累加。Tensor Core的思路完全不同既然这些操作模式固定、计算密集、不需要复杂控制流那我直接在硬件里塞专门的矩阵乘法单元一次指令直接算一个4x4矩阵的乘累加。A100的Tensor Core每个时钟周期可以完成一个4x4x4的矩阵运算一个SM里还有多个Tensor Core同时工作。数据上看A100的FP32算力是19.5 TFLOPS但Tensor Core的FP16算力高达312 TFLOPS含稀疏差了整整16倍。也就是说同样面积的硅片跑通用的CUDA Core指令只能达到这个速度切到Tensor Core后算力直接上了一个数量级。3.2 精度控制FP16、BF16与FP8Tensor Core刚推出时默认搭配FP16计算。FP16的问题在于动态范围小数值容易溢出或下溢。训练过程中梯度一波动FP16可能就爆了。对策是混合精度训练也就是主权重和梯度保存在FP32计算时把数据cast成FP16梯度累加回FP32同时用loss scaling保证梯度不爆炸。这套方案已经是训练标配了。到了Ampere架构NVIDIA又加入了BF16bfloat16。BF16牺牲了尾数位换来了与FP32相同的指数范围让混合精度训练的稳定性大幅提升不再需要那么繁琐的loss scaling策略。也是在这代架构上Tensor Core算力从V100的125 TFLOPS涨到A100的312 TFLOPS。Hopper架构加入FP8支持。FP8的好处是进一步减半数据量Tensor Core吞吐理论上再加倍。H100的FP8稀疏算力是1979 TFLOPS这个数字看着夸张但对精度极其敏感现在主要用在推理场景或已经被验证稳健的训练任务里比如FP8版DeepSeek-R1的训练或vLLM的FP8推理。3.3 Tensor Core使用门槛需要框架层配合Tensor Core不是写普通CUDA代码就能自动用上的而是需要显式调用cuBLAS、cuDNN这类库或者你在CUDA C里用WMMAwarp matrix multiply-accumulate指令写PTX更贴近硬件的用mma.sync指令直接搓。大部分AI工程师不需要到这个层面因为PyTorch用的cuBLAS已经帮你把路径调好了。但如果你是做自定义算子、量化推理、或者写高性能kernel的就必须知道什么数据排布比如按K维度切块、用ldmatrix指令加载到共享内存才能喂饱Tensor Core——多数性能不达标的自定义算子问题就出在这里。重要的认知Tensor Core的算力虽然暴力但它吃数据的能力更强。一个Tensor Core指令周期极短如果你没有把数据从显存搬到寄存器再喂给它它就在空转。这也是为什么AI计算里内存带宽往往比峰值算力更值得关注。4. 显存带宽与互联算力以外真正的瓶颈很多人的性能分析路径是先看FLOPs再看利用率最后发现利用率低然后把锅甩给框架。但这个链条的起点其实是错的。真实情况往往是你根本没有喂饱算力因为在数据搬运环节已经卡脖子了。4.1 从GPU内部看内存层次从寄存器到L1/共享内存、L2缓存再到HBM显存每一级的带宽逐级下降延迟逐级上升。寄存器的访问延迟是纳秒级带宽是TB/s级别到了HBM这一层A100的带宽是2TB/sH100是3.35TB/s看着很大但对比算力需求就捉襟见肘了。一个简单的算术假设我们要做FP16矩阵乘法H100的Tensor Core算力是989 TFLOPS不含稀疏。每个FP16数2字节一次乘加需要读两个数乘数A和B那就是4字节的输入。989 TFLOPS意味着每秒钟需要989万亿次乘加对应需要大约4PB/s的数据供给989万亿 x 4字节。但H100的HBM带宽只有3.35TB/s相差了三个数量级。也就是说如果数据全部从显存读取H100的算力利用率连1%都到不了。这就是为什么矩阵乘法一定需要分块计算把数据块加载到共享内存/寄存器里反复复用让每个字节从HBM读进来以后参加多次计算把这个AI算力 / 显存带宽的比值业界叫算术强度arithmetic intensity压下去性能才能起来。4.2 多卡互联NVLink与PCIe的差距单卡的问题解决了多卡通信的带宽更是从物理层面限制大模型训练的关键。PCIe 4.0 x16的单向带宽是32GB/s双向64GB/s。听起来也不低了。但NVIDIA自家GPU之间走的是NVLinkA100的NVLink 3.0是600GB/s双向H100的NVLink 4.0达到900GB/s最新Blackwell平台的NVLink 5.0有1800GB/s。一个数量级的差距这就是为什么训练大模型时推荐尽量走NVLink交换网络而不是PCIe。做AI Infra的人得明白这个数差距意味着什么。假设一个batch的数据量是100GB用PCIe传输需要3秒用NVLink只需0.1秒出头。在大规模训练里这个差异直接决定了你在通信上花费的时间占比。4.3 NPS/MFU理解真实算力利用率的钥匙前面说了那么多架构知识最终落地到两个指标一个叫MFUModel FLOPs Utilization模型浮点运算利用率一个叫NPSNormalized Performance Score归一化性能分数。MFU定义为实际算力除以理论峰值算力。比如训练一个7B模型H100峰值摩尔989 TFLOPS实测整个训练过程平均有效算力400 TFLOPSMFU就是40%。这不是一个简单数字它已经综合了数据加载、kernel效率、通信开销、气泡等待等所有损耗。上一届A100训练的大模型MFU普遍在35%-45%之间优化到50%以上就是非常优秀的团队。H100因为算力和架构双双提升好的训练配置能把MFU抬到55%-60%。NPS是MLPerf里用的归一化指标把不同加速器的实际吞吐除以各自的理论峰值也得到一个百分比。有了这两个概念你就不会看到占满GPU就以为算力用满了而会去追问——占的是算力还是带宽是计算单元在忙还是搬运数据在忙。5. 架构演进路线图从Ampere到Hopper到Blackwell理解了架构底层以后去看看最近的几代架构是怎么演进的你会更清楚NVIDIA在想什么。5.1 AmpereA100AI训练的分水岭A100发布于2020年基于GA100芯片。它第一次把Tensor Core的算力拉到一个恐怖的量级同时加入第三代NVLink600GB/s和40GB HBM2e显存。对这代架构的记忆主要有两个一是它让混合精度训练成为绝对主流二是引入了MIGMulti-Instance GPU可以把一张卡切成多个实例。这个切分能力对AI Infra的集群调度很有价值可以让多任务共享一张大卡。5.2 HopperH100/H200Transformer专用化Hopper相比Ampere最核心的变化是加入了Transformer Engine和FP8支持。这代架构是深度绑定Transformer模型的——通过动态地在FP8/FP16之间切换精度让Transformer的forward/backward跑得更快。另一个重要变化是引入TMATensor Memory Accelerator。老GPU搬运数据要靠线程显式发起load/storeTMA是一个独立的硬件单元可以直接把显存数据搬到共享内存不需要线程参与搬运。线程等待数据的时间大幅减少本质上是用硬件帮软件程序员解决了一部分内存带宽利用问题。H100的SM规模比A100多出22%共享内存容量也翻倍这都直接利好kernel性能。后续的H200在显存上从80GB升到141GBHBM3e带宽从3.35TB/s涨到4.8TB/s——纯粹为大模型参数更大地驻留在单卡上做的升级。5.3 BlackwellAI Infra视角的又一次跳跃Blackwell B200把两个die封装在一起用高带宽die-to-die互联单卡看似是两块GPU但逻辑上是一个设备。它引入了FP4支持NVLink 5.0带宽大幅提升到1800GB/s显存从141GB翻到288GB。这些都不意外——继续增加单卡算力和显存容量。真正值得AI Infra团队注意的是它的形态如果你要搭集群你买的不再是卡更像是一台台GPU服务器节点。风冷和液冷的选择从B200开始分化H100时代风冷还勉强压得住B200这种级别的功耗单卡TDP超过1000W基本只剩液冷一条路。集群设计里的供电、散热、机柜空间规划都要重新核算。5.4 国内GPU生态的架构差异顺带提一嘴国产GPU。华为昇腾的Ascend 910B3、寒武纪的思元590、算能的SG2300系列都已经能跑大模型训练。昇腾910B3对标A100级别的算力910C接近H100的算力水平。不过生态差异仍然不容忽视——CUDA的软件栈、PyTorch的原生支持、各种开源库的适配国产GPU暂时还达不到同样的顺滑程度。从AI Infra的架构设计角度如果要用国产卡你得在算子适配、框架迁移、通信库替换比如华为的HCCL对标NCCL这些层面多准备不少工作量还要接受性能调优上的各种坑。这也是为什么很多公司做选型时即便国产卡价格低且供货稳定依然要把CUDA生态作为重要考量维度——这不是偏见这是现实。6. 从入门到实践如何用实测数据反推瓶颈纸上谈兵聊完架构最重要的还是回到工作场景你有一台GPU机器训练跑不快怎么定位到底是哪里不行6.1 测量第一Nsight Systems和Nsight Compute先不要猜先测。NVIDIA提供两个官方工具它们的使用场景完全不同。Nsight Systems是系统级的看CPU和GPU整体时间线kernel有没有密集执行CPU有没有提前把数据准备好通信和计算有没有重叠。经常能发现的问题是GPU在等CPU——CPU下发指令太慢GPU空转。这种情况下的优化方向是加大batch size、用DataLoader多线程预取、或者减少Python端同步等待。Nsight Compute是kernel级的看具体kernel内部的执行效率GPU吞吐利用率是多少瓶颈在load/store还是在计算寄存器用了多少有没有bank conflictwarp有没有idle。这把工具可以看到每个kernel的enqueue到执行的全链路指标。6.2 用Profiling结果对照架构模型举个例子我在调一个批量矩阵乘法场景时发现Nsight Compute报告显示warp stall原因一栏long scoreboard占90%——这个指标的意思是线程在等待数据从显存返回。也就是说kernel不是算不动而是数据喂不过来。优化方向立刻就清楚了增加每个线程的数据复用、用共享内存做block级缓存、或者调整数据排布让L2命中率提高。如果报告显示math pipe throttle那说明计算管线已经满了也就是每周期都有指令要计算你接近算力上限。这时候再去改数据复用已无济于事性价比更高的方向是切换到更省算力的低精度模式。6.3 从GPU架构视角看框架层优化项知道架构再看框架层的优化项工作就变得立体了张量内存分配PyTorch默认的缓存分配器是把显存踩得更碎还是更整取决于你的tensor生命周期管理。用好torch.cuda.memory的snapshot确认没有大量碎片化空闲块。数据排布深度学习中tensor默认是NCHW通道在前但GPU上很多算子实际期望NHWC或特定对齐格式。改变内存布局往往能显著提升Tensor Core利用率尤其是CNN场景。算子融合多个小kernel连续执行每完成一个都要回显存。把这几个算子合并在一个kernel里中间结果留在寄存器或共享内存省掉多次显存往返。这是PyTorch 2.0的torch.compile和CUDA graph在干的事情之一。这些策略的底层动机统统指向一个架构事实GPU的算力在飞速增长但数据搬运的代价没人帮你免掉框架层优化本质上是绕过数据搬运瓶颈。7. GPU架构知识在AI Infra实践中的延伸把架构知识落地到日常的AI Infra工作中还有几个跟架构直接相关的选择问题值得展开。7.1 卡型选型和集群规划的关系做AI Infra经常要做卡型选型。涉及GPU架构的知识时不要只看单卡算力要看你负载的形态。跑大模型预训练显存容量和带宽权重极高。模型参数、优化器状态、梯度、KV Cache全部要驻留在显存里单卡80GB放不小得张量并行加流水线并行拆到多卡。所以H200/A100 80G这种大显存卡是训练集群的基本选择。做推理服务关键指标变成响应延迟和吞吐这时FP8/FP4时代的小显存卡、甚至像L40S那样价格折中的卡性价比可能更高。AI Infra的人如果选错卡差的不是一星半点的成本而是几百万的设备预算被浪费。7.2 显存容量翻倍对并行策略的直接影响H200的141GB对比H100的80GB不是简单的参数加多深度学习并行策略会因此改变。一个7B模型用BF16训练参数权重加梯度加优化器状态大约需要7GB乘以若干倍单卡能装下意味着可以免去部分张量并行。减少张量并行意味着跨卡通信量下降训练吞吐自然提升。所以显存容量和MFU之间存在隐形的联动关系架构上理解不到这一层做集群设计时就只能拍脑袋。7.3 液冷与供电架构功耗带来的现实问题上面提到B200功耗破千W国内很多机房目前的机柜功率密度不够强行上液冷你还要做CDU选型、漏液检测、冷却液回路规划。GPU架构演进推高了单节点功率密度AI Infra的活不只是停留在软件层硬件设施甚至机房选址建设同样要被纳入规划。8. 踩坑记录调优路上最常误判的几个场景最后分享几个我踩过的坑都是真实项目中因为对GPU架构理解不到位导致的误判。第一个坑是盲目升级到更大算力的卡。曾经有个服务场景推理的batch size一直提不上去以为是算力不够换了大卡以后吞吐几乎没变后来用Nsight看了一下瓶颈全在数据加载路径上——CPU侧的数据预处理跟不上GPU百分之九十时间在等。换了卡是白花钱。第二个坑是只看HBM带宽买卡。当时比较两张卡一个带宽高一些但L2缓存小一个带宽低一点但L2缓存大最终选了高带宽的。实际负载中因为数据复用率高L2命中率大才是更重要的倒过来选反而更划算。所以选卡不能只看纸面参数要结合自己的负载形态去看profile数据。第三个坑是轻视PCIe带宽。搭一台多卡推理机时贪便宜用了PCIe交换而不是NVLink发现模型并行时通信瓶颈明显MFU掉了近二十个点。做集群方案时通信互联的预算不能省。第四个坑更隐蔽共享内存和bank conflict。写自定义attention kernel时共享内存命中率很高但带宽就是上不去。后来查出来是bank conflict导致同一地址被多个线程同时访问串行化了。这个在Nsight Compute的指标里清清楚楚但你要看得懂shared bank conflict这个指标。这些坑都在提醒我同一个道理GPU不是一个黑盒它是一台有明确取舍规律的并行机器。会用它的编程框架只是第一层理解底层的吞吐逻辑你才算真正入了AI Infra的门。如果你想更深入地学建议花点时间看NVIDIA的CUDA C Programming Guide里的编程模型章节配合Nsight系列工具做一遍profile练习。理论模型和实测数据一一对上之后那种整个架构在脑子里亮起来的感觉是后面所有优化工作的底气。
