SemiAnalysis Massive Teardown: Full Blackwell Architecture Details, NVIDIA's Never-Before-Revealed Secrets

英伟达 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 个百分点。