SemiAnalysis首度拆解英伟达Blackwell架构：在AI负载下，张量核心与内存带宽整体逼近理论峰值，但性能高度依赖指令形状与软件调优。2SM MMA实现近乎完美扩展，但SMEM带宽与跨Die约300周期延迟成为关键瓶颈。研究揭示，Blackwell性能释放不取决于硬件上限，而取决于调度与优化能力。
英伟达Blackwell GPU代表了近年来最重大的GPU微架构变革之一，但迄今缺乏详尽的官方白皮书。
知名半导体研究机构SemiAnalysis历时数月，对Blackwell架构进行了系统性微基准测试，首次公开了该架构在AI工作负载下的硬件性能上限数据。
测试结果显示，Blackwell在张量核心（Tensor Core）吞吐量、内存子系统带宽及新型2SM MMA指令等关键维度上均接近理论峰值，但性能表现高度依赖指令形状配置，部分场景下存在明显的带宽瓶颈。这一发现对AI基础设施投资者和芯片采购方具有直接参考价值——架构潜力能否充分释放，取决于软件层面的精细调优。
SemiAnalysis已将相关基准测试代码库开源，测试所用B200节点由Nebius和Verda提供。研究团队同时宣布，后续将扩展至TPU Pallas内核、Trainium NKI内核及AMD CDNA4汇编的基准测试。
架构核心变化：TMEM引入与2SM MMA
从Hopper到Blackwell，英伟达对MMA相关指令的PTX抽象层进行了多项重要调整。
最显著的变化是引入了张量内存（TMEM）用于存储MMA累加器。在此前架构中，线程隐式持有MMA运算结果；Blackwell改为由软件在MMA作用域内显式管理TMEM，改变了线程与计算结果之间的所有权关系。
与此同时，tcgen05操作现在由单一线程代表整个CTA（协作线程阵列）发出，而非此前Hopper架构中以warp或warpgroup为单位发出。这一变化在CuTe MMA原子中有直接体现：Blackwell使用ThrID = Layout<_1>，而Hopper使用ThrID = Layout<_128>。
Blackwell还引入了TPC作用域的TMA和MMA，支持两个协同CTA跨SM对执行tcgen05.mma，共享操作数，从而在降低每个CTA共享内存带宽需求的同时，提供更高运算强度的MMA指令。此外，该架构原生支持带微缩放的亚字节数据类型，并引入了集群启动控制（CLC）作为持久化CTA内核中动态工作调度的硬件支持。
芯片物理布局：双Die架构与300周期跨Die延迟
SemiAnalysis通过逆向工程手段，揭示了B200芯片的物理拓扑结构。
研究团队利用PTX %%smid指令，通过启动不同大小的集群来反向推断SM到GPC（图形处理集群）的映射关系。结果显示，B200存在部分TPC独占逻辑GPC的情况，这些TPC从不与其他TPC协同调度。
通过让每个SM遍历填满L2缓存的指针追踪数组并测量各SM间的访问延迟，研究团队构建了SM间距离矩阵。矩阵清晰呈现出两组SM，平均L2访问延迟差距超过300个时钟周期，对应的正是两个Die之间的跨Die访问惩罚。
基于此，研究团队推断B200的Die级TPC分布如下：
Die A：各GPC分别包含10、10、10、9个TPC
Die B：各GPC分别包含9、9、9、5+3个TPC
这一物理布局差异意味着，即便逻辑配置相同的两块GPU，其物理SM分布也可能不同，构成潜在的性能非确定性来源。
内存子系统：LDGSTS与TMA的性能边界
内存子系统测试聚焦于两类异步拷贝指令：LDGSTS（异步拷贝）和TMA（张量内存加速器）。
LDGSTS方面，测试覆盖了FlashInfer多头注意力（MHA）内核的典型配置。结果显示，LDGSTS内存吞吐量在32 KiB在途字节时饱和，峰值约为6.6 TB/s。16字节加载在相同在途字节数下略优于8字节加载，且消耗更少执行资源。延迟测试显示，LDGSTS基线延迟约为600纳秒，在途字节超过8 KiB后延迟接近翻倍，原因在于大量线程因MIO（内存输入输出）节流而停滞。
TMA方面，峰值吞吐量的达到明显晚于LDGSTS。在低于32字节在途数据时，异步拷贝吞吐量略优于TMA；超过该阈值后TMA追上并可持续扩展至128 KiB。延迟方面，在途数据低于12 KiB时异步拷贝延迟略低，超过后TMA延迟大幅攀升。
TMA多播测试显示，显式TMA多播可完美消除L2流量，实现理想的"1/集群大小"L2字节比。隐式多播（各CTA独立发出TMA加载至相同数据）在有效内存吞吐量上与显式多播相当，但在超过64字节在途数据后，L2缓存流量削减效果开始下降。
张量核心性能：形状依赖性显著，2SM MMA实现完美弱扩展
张量核心测试是本次研究的核心部分，结果揭示了Blackwell MMA性能对指令形状的高度敏感性。
吞吐量方面，对于1SM MMA，M=64的配置最高仅能达到理论峰值的50%，而M=128可接近100%。这证实M=64仅利用了一半数据通路。对于2SM MMA，M=128在N=64时吞吐量为峰值的90%，其余N尺寸均接近100%；M=256则在所有配置下均维持接近100%的峰值吞吐量，因为M=256等效于每SM处理M=128，可充分利用完整数据通路。
AB布局影响同样显著。当两个输入矩阵均存储于共享内存（SS模式）时，M=128在N<128时存在明显的SMEM带宽瓶颈。以FP16为例，硬件每周期可执行8192 MMA FLOP，SMEM带宽为128 B/周期，计算表明M=128 N=64 K=16配置下SMEM需要48个周期，而数学运算仅需32个周期，即指令受SMEM带宽限制。所有数据类型均存在这一规律——双操作数均在SMEM中的MMA指令，在N<128时均受SMEM带宽约束。
2SM MMA实现了完美的弱扩展，相对于1SM MMA在使用两倍计算资源时获得2倍加速。在SS模式的小形状配置下，由于操作数B在两个SM间分片，甚至出现超过2倍的加速。研究结论明确：应始终使用给定SMEM tile尺寸下可用的最大指令形状，以获得最高吞吐量。
延迟方面，所有配置下延迟均随N从64增至128线性增长，N=256时出现跳跃。数据类型延迟排序呈现规律性：S8 < BF16 = E4M3 = F4 < MXF8 = MXF4，研究团队认为整数运算功耗效率更高导致S8最快，而微缩放数据类型的缩放因子计算引入了轻微额外开销。
实际在途指令数测试显示，在典型内核使用的1至4条在途MMA指令场景下，4条在途MMA的吞吐量上限约为理论峰值的78%至80%，且1SM MMA比2SM MMA高出约5个百分点。
Comments