文章核心观点 - 英伟达Blackwell GPU架构在张量核心吞吐量、内存带宽等关键硬件指标上接近理论峰值，但性能表现高度依赖指令形状配置，部分场景存在带宽瓶颈，其潜力释放取决于软件层面的精细调优 [1] 架构核心变化 - 引入了张量内存（TMEM）用于显式管理MMA累加器，改变了线程与计算结果的所有权关系 [2] - 引入了tcgen05操作由单一线程代表整个CTA发出，而Hopper架构以warp或warpgroup为单位发出 [2] - 引入了TPC作用域的TMA和MMA，支持两个协同CTA跨SM执行tcgen05.mma，共享操作数以降低每个CTA的共享内存带宽需求 [2] - 原生支持带微缩放的亚字节数据类型，并引入了集群启动控制（CLC）作为硬件支持 [2] 芯片物理布局 - B200芯片采用双Die架构，两组SM之间的平均L2访问延迟差距超过300个时钟周期，对应跨Die访问惩罚 [3] - Die A的GPC分别包含10、10、10、9个TPC，Die B的GPC分别包含9、9、9、5+3个TPC [4] - 物理布局差异意味着逻辑配置相同的GPU，其物理SM分布也可能不同，构成潜在的性能非确定性来源 [3] 内存子系统性能 - LDGSTS（异步拷贝）内存吞吐量在32 KiB在途字节时饱和，峰值约为6.6 TB/s [5] - 16字节加载在相同在途字节数下略优于8字节加载，且消耗更少执行资源 [5] - LDGSTS基线延迟约为600纳秒，在途字节超过8 KiB后延迟接近翻倍 [5] - TMA（张量内存加速器）在低于32字节在途数据时，吞吐量略低于异步拷贝；超过该阈值后TMA追上并可持续扩展至128 KiB [6] - 在途数据低于12 KiB时异步拷贝延迟略低，超过后TMA延迟大幅攀升 [6] - 显式TMA多播可完美消除L2流量，实现理想的"1/集群大小"L2字节比 [6] - 隐式多播在有效内存吞吐量上与显式多播相当，但在超过64字节在途数据后，L2缓存流量削减效果下降 [6] 张量核心性能 - 张量核心性能对指令形状高度敏感 [7] - 对于1SM MMA，M=64的配置最高仅能达到理论峰值的50%，而M=128可接近100% [7] - 对于2SM MMA，M=128在N=64时吞吐量为峰值的90%，其余N尺寸均接近100%；M=256在所有配置下均维持接近100%的峰值吞吐量 [7] - 当两个输入矩阵均存储于共享内存（SS模式）时，M=128在N<128时存在明显的SMEM带宽瓶颈 [7] - 2SM MMA实现了完美的弱扩展，相对于1SM MMA在使用两倍计算资源时获得2倍加速，在SS模式的小形状配置下甚至出现超过2倍的加速 [8] - 所有配置下延迟均随N从64增至128线性增长，N=256时出现跳跃 [8] - 数据类型延迟排序为：S8 < BF16 = E4M3 = F4 < MXF8 = MXF4 [8] - 在1至4条在途MMA指令的典型场景下，4条在途MMA的吞吐量上限约为理论峰值的78%至80%，且1SM MMA比2SM MMA高出约5个百分点 [9]