最近这阵子团队准备自建训练集群，方案评审时经常被问到同一个问题：为什么花这么多钱买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则是SIMT（Single 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 SM：GPU里的"计算集群"
NVIDIA GPU里最核心的计算单元叫SM（Streaming Multiprocessor，流式多处理器）。你可以把SM理解成一个小型的计算集群，里面有自己的执行单元、寄存器文件、共享内存和调度器。一个GPU里通常有几十到上百个SM，比如A100有108个SM，H100有132个SM，H800同样132个SM。
每个SM内部又包含若干CUDA Core（确切地说是FP32单元，不过是通用叫法）、Tensor Core（张量核心，后文细讲）、特殊函数单元（SFU），以及负责调度这些单元的调度器。
看一台GPU的算力，表面上是看FP32 TFLOPS这个数字，本质上是看SM的数量乘以每个SM每周期能执行的运算量。A100的FP32算力19.5 TFLOPS，H100是67 TFLOPS（带稀疏性），数字差距背后就是SM规模和每周期执行能力的双重差异。
2.2 Warp：GPU执行调度的基本单位
这是理解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/64=1024个线程。如果寄存器不够用，会发生什么？变量会溢出（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的乘累加操作（FMA，fused 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与FP8
Tensor Core刚推出时，默认搭配FP16计算。FP16的问题在于动态范围小，数值容易溢出或下溢。训练过程中梯度一波动，FP16可能就爆了。
对策是混合精度训练，也就是主权重和梯度保存在FP32，计算时把数据cast成FP16，梯度累加回FP32，同时用loss scaling保证梯度不爆炸。这套方案已经是训练标配了。
到了Ampere架构，NVIDIA又加入了BF16（bfloat16）。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++里用WMMA（warp 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/s，H100是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之间走的是NVLink，A100的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：理解真实算力利用率的钥匙
前面说了那么多架构知识，最终落地到两个指标：一个叫MFU（Model FLOPs Utilization，模型浮点运算利用率），一个叫NPS（Normalized Performance Score，归一化性能分数）。
MFU定义为实际算力除以理论峰值算力。比如训练一个7B模型，H100峰值摩尔989 TFLOPS，实测整个训练过程平均有效算力400 TFLOPS，MFU就是40%。这不是一个简单数字，它已经综合了数据加载、kernel效率、通信开销、气泡等待等所有损耗。
上一届A100训练的大模型MFU普遍在35%-45%之间，优化到50%以上就是非常优秀的团队。H100因为算力和架构双双提升，好的训练配置能把MFU抬到55%-60%。
NPS是MLPerf里用的归一化指标，把不同加速器的实际吞吐除以各自的理论峰值，也得到一个百分比。有了这两个概念，你就不会看到"占满GPU"就以为算力用满了，而会去追问——占的是算力还是带宽？是计算单元在忙还是搬运数据在忙。
5. 架构演进路线图：从Ampere到Hopper到Blackwell
理解了架构底层以后，去看看最近的几代架构是怎么演进的，你会更清楚NVIDIA在想什么。
5.1 Ampere（A100）：AI训练的分水岭
A100发布于2020年，基于GA100芯片。它第一次把Tensor Core的算力拉到一个恐怖的量级，同时加入第三代NVLink（600GB/s）和40GB HBM2e显存。对这代架构的记忆主要有两个：一是它让混合精度训练成为绝对主流，二是引入了MIG（Multi-Instance GPU）可以把一张卡切成多个实例。这个切分能力对AI Infra的集群调度很有价值，可以让多任务共享一张大卡。
5.2 Hopper（H100/H200）：Transformer专用化
Hopper相比Ampere，最核心的变化是加入了Transformer Engine和FP8支持。这代架构是深度绑定Transformer模型的——通过动态地在FP8/FP16之间切换精度，让Transformer的forward/backward跑得更快。
另一个重要变化是引入TMA（Tensor Memory Accelerator）。老GPU搬运数据要靠线程显式发起load/store，TMA是一个独立的硬件单元，可以直接把显存数据搬到共享内存，不需要线程参与搬运。线程等待数据的时间大幅减少，本质上是用硬件帮软件程序员解决了一部分内存带宽利用问题。
H100的SM规模比A100多出22%，共享内存容量也翻倍，这都直接利好kernel性能。后续的H200在显存上从80GB升到141GB（HBM3e），带宽从3.35TB/s涨到4.8TB/s——纯粹为大模型参数更大地驻留在单卡上做的升级。
5.3 Blackwell：AI 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 conflict，warp有没有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练习。理论模型和实测数据一一对上之后，那种"整个架构在脑子里亮起来"的感觉，是后面所有优化工作的底气。
