Blackwell架构核心观点 - 知名半导体研究机构SemiAnalysis通过系统性微基准测试，首次公开了Blackwell架构在AI工作负载下的硬件性能上限数据[3] - 测试结果显示，Blackwell在张量核心吞吐量、内存子系统带宽及新型2SM MMA指令等关键维度上均接近理论峰值，但性能表现高度依赖指令形状配置，部分场景下存在明显的带宽瓶颈[3] - 研究揭示，Blackwell性能释放不取决于硬件上限，而取决于软件层面的精细调优与调度能力[2][3] 架构核心变化 - 从Hopper到Blackwell，英伟达引入了张量内存（TMEM）用于显式管理MMA累加器，改变了线程与计算结果之间的所有权关系[6] - tcgen05操作现在由单一线程代表整个CTA发出，而Hopper架构中以warp或warpgroup为单位发出[6] - 架构引入了TPC作用域的TMA和MMA，支持两个协同CTA跨SM对执行tcgen05.mma，共享操作数，从而在降低每个CTA共享内存带宽需求的同时，提供更高运算强度[6] - 架构原生支持带微缩放的亚字节数据类型，并引入了集群启动控制（CLC）作为持久化CTA内核中动态工作调度的硬件支持[6] 芯片物理布局 - SemiAnalysis通过逆向工程揭示B200芯片采用双Die物理拓扑结构，两组SM之间的平均L2访问延迟差距超过300个时钟周期，对应两个Die之间的跨Die访问惩罚[8][9] - 研究推断B200的Die级TPC分布不均，Die A各GPC分别包含10、10、10、9个TPC，Die B各GPC分别包含9、9、9、5+3个TPC[10] - 物理布局差异意味着，即便逻辑配置相同的两块GPU，其物理SM分布也可能不同，构成潜在的性能非确定性来源[10] 内存子系统性能 - 异步拷贝指令LDGSTS的内存吞吐量在32 KiB在途字节时饱和，峰值约为6.6 TB/s，16字节加载在相同在途字节数下略优于8字节加载[12] - LDGSTS基线延迟约为600纳秒，在途字节超过8 KiB后延迟接近翻倍，原因在于大量线程因MIO节流而停滞[13] - 显式TMA多播可完美消除L2流量，实现理想的“1/集群大小”L2字节比，隐式多播在有效内存吞吐量上与之相当，但在超过64字节在途数据后，L2缓存流量削减效果开始下降[14] - TMA达到峰值吞吐量明显晚于LDGSTS，在低于32字节在途数据时异步拷贝吞吐量略优于TMA，超过该阈值后TMA追上并可持续扩展至128 KiB[15] 张量核心性能 - Blackwell MMA性能对指令形状高度敏感，对于1SM MMA，M=64的配置最高仅能达到理论峰值的50%，而M=128可接近100%[18] - 对于2SM MMA，M=128在N=64时吞吐量为峰值的90%，其余N尺寸均接近100%，M=256则在所有配置下均维持接近100%的峰值吞吐量[18] - 当两个输入矩阵均存储于共享内存（SS模式）时，M=128在N<128时存在明显的SMEM带宽瓶颈，所有数据类型均存在这一规律[20] - 2SM MMA实现了完美的弱扩展，相对于1SM MMA在使用两倍计算资源时获得2倍加速，在SS模式的小形状配置下，由于操作数B在两个SM间分片，甚至出现超过2倍的加速[20] - 研究结论明确：应始终使用给定SMEM tile尺寸下可用的最大指令形状，以获得最高吞吐量[20] - 在典型内核使用的1至4条在途MMA指令场景下，4条在途MMA的吞吐量上限约为理论峰值的78%至80%，且1SM MMA比2SM MMA高出约5个百分点[22]