英偉達(dá)Blackwell GPU代表了近年來最重大的GPU微架構(gòu)變革之一，但迄今缺乏詳盡的官方白皮書。

知名半導(dǎo)體研究機(jī)構(gòu)SemiAnalysis歷時(shí)數(shù)月，對Blackwell架構(gòu)進(jìn)行了系統(tǒng)性微基準(zhǔn)測試，首次公開了該架構(gòu)在AI工作負(fù)載下的硬件性能上限數(shù)據(jù)。

測試結(jié)果顯示，Blackwell在張量核心（Tensor Core）吞吐量、內(nèi)存子系統(tǒng)帶寬及新型2SM MMA指令等關(guān)鍵維度上均接近理論峰值，但性能表現(xiàn)高度依賴指令形狀配置，部分場景下存在明顯的帶寬瓶頸。這一發(fā)現(xiàn)對AI基礎(chǔ)設(shè)施投資者和芯片采購方具有直接參考價(jià)值——架構(gòu)潛力能否充分釋放，取決于軟件層面的精細(xì)調(diào)優(yōu)。

SemiAnalysis已將相關(guān)基準(zhǔn)測試代碼庫開源，測試所用B200節(jié)點(diǎn)由Nebius和Verda提供。研究團(tuán)隊(duì)同時(shí)宣布，后續(xù)將擴(kuò)展至TPU Pallas內(nèi)核、Trainium NKI內(nèi)核及AMD CDNA4匯編的基準(zhǔn)測試。

架構(gòu)核心變化：TMEM引入與2SM MMA

從Hopper到Blackwell，英偉達(dá)對MMA相關(guān)指令的PTX抽象層進(jìn)行了多項(xiàng)重要調(diào)整。

最顯著的變化是引入了張量內(nèi)存（TMEM）用于存儲MMA累加器。在此前架構(gòu)中，線程隱式持有MMA運(yùn)算結(jié)果；Blackwell改為由軟件在MMA作用域內(nèi)顯式管理TMEM，改變了線程與計(jì)算結(jié)果之間的所有權(quán)關(guān)系。

與此同時(shí)，tcgen05操作現(xiàn)在由單一線程代表整個(gè)CTA（協(xié)作線程陣列）發(fā)出，而非此前Hopper架構(gòu)中以warp或warpgroup為單位發(fā)出。這一變化在CuTe MMA原子中有直接體現(xiàn)：Blackwell使用ThrID = Layout<_1>，而Hopper使用ThrID = Layout<_128>。

Blackwell還引入了TPC作用域的TMA和MMA，支持兩個(gè)協(xié)同CTA跨SM對執(zhí)行tcgen05.mma，共享操作數(shù)，從而在降低每個(gè)CTA共享內(nèi)存帶寬需求的同時(shí)，提供更高運(yùn)算強(qiáng)度的MMA指令。此外，該架構(gòu)原生支持帶微縮放的亞字節(jié)數(shù)據(jù)類型，并引入了集群啟動控制（CLC）作為持久化CTA內(nèi)核中動態(tài)工作調(diào)度的硬件支持。

芯片物理布局：雙Die架構(gòu)與300周期跨Die延遲

SemiAnalysis通過逆向工程手段，揭示了B200芯片的物理拓?fù)浣Y(jié)構(gòu)。

研究團(tuán)隊(duì)利用PTX %%smid指令，通過啟動不同大小的集群來反向推斷SM到GPC（圖形處理集群）的映射關(guān)系。結(jié)果顯示，B200存在部分TPC獨(dú)占邏輯GPC的情況，這些TPC從不與其他TPC協(xié)同調(diào)度。

通過讓每個(gè)SM遍歷填滿L2緩存的指針追蹤數(shù)組并測量各SM間的訪問延遲，研究團(tuán)隊(duì)構(gòu)建了SM間距離矩陣。矩陣清晰呈現(xiàn)出兩組SM，平均L2訪問延遲差距超過300個(gè)時(shí)鐘周期，對應(yīng)的正是兩個(gè)Die之間的跨Die訪問懲罰。

基于此，研究團(tuán)隊(duì)推斷B200的Die級TPC分布如下：

• Die A：各GPC分別包含10、10、10、9個(gè)TPC

• Die B：各GPC分別包含9、9、9、5+3個(gè)TPC

這一物理布局差異意味著，即便邏輯配置相同的兩塊GPU，其物理SM分布也可能不同，構(gòu)成潛在的性能非確定性來源。

內(nèi)存子系統(tǒng)：LDGSTS與TMA的性能邊界

內(nèi)存子系統(tǒng)測試聚焦于兩類異步拷貝指令：LDGSTS（異步拷貝）和TMA（張量內(nèi)存加速器）。

LDGSTS方面，測試覆蓋了FlashInfer多頭注意力（MHA）內(nèi)核的典型配置。結(jié)果顯示，LDGSTS內(nèi)存吞吐量在32 KiB在途字節(jié)時(shí)飽和，峰值約為6.6 TB/s。16字節(jié)加載在相同在途字節(jié)數(shù)下略優(yōu)于8字節(jié)加載，且消耗更少執(zhí)行資源。延遲測試顯示，LDGSTS基線延遲約為600納秒，在途字節(jié)超過8 KiB后延遲接近翻倍，原因在于大量線程因MIO（內(nèi)存輸入輸出）節(jié)流而停滯。

TMA方面，峰值吞吐量的達(dá)到明顯晚于LDGSTS。在低于32字節(jié)在途數(shù)據(jù)時(shí)，異步拷貝吞吐量略優(yōu)于TMA；超過該閾值后TMA追上并可持續(xù)擴(kuò)展至128 KiB。延遲方面，在途數(shù)據(jù)低于12 KiB時(shí)異步拷貝延遲略低，超過后TMA延遲大幅攀升。

TMA多播測試顯示，顯式TMA多播可完美消除L2流量，實(shí)現(xiàn)理想的"1/集群大小"L2字節(jié)比。隱式多播（各CTA獨(dú)立發(fā)出TMA加載至相同數(shù)據(jù)）在有效內(nèi)存吞吐量上與顯式多播相當(dāng)，但在超過64字節(jié)在途數(shù)據(jù)后，L2緩存流量削減效果開始下降。

張量核心性能：形狀依賴性顯著，2SM MMA實(shí)現(xiàn)完美弱擴(kuò)展

張量核心測試是本次研究的核心部分，結(jié)果揭示了Blackwell MMA性能對指令形狀的高度敏感性。

吞吐量方面，對于1SM MMA，M=64的配置最高僅能達(dá)到理論峰值的50%，而M=128可接近100%。這證實(shí)M=64僅利用了一半數(shù)據(jù)通路。對于2SM MMA，M=128在N=64時(shí)吞吐量為峰值的90%，其余N尺寸均接近100%；M=256則在所有配置下均維持接近100%的峰值吞吐量，因?yàn)镸=256等效于每SM處理M=128，可充分利用完整數(shù)據(jù)通路。

AB布局影響同樣顯著。當(dāng)兩個(gè)輸入矩陣均存儲于共享內(nèi)存（SS模式）時(shí)，M=128在N<128時(shí)存在明顯的SMEM帶寬瓶頸。以FP16為例，硬件每周期可執(zhí)行8192 MMA FLOP，SMEM帶寬為128 B/周期，計(jì)算表明M=128 N=64 K=16配置下SMEM需要48個(gè)周期，而數(shù)學(xué)運(yùn)算僅需32個(gè)周期，即指令受SMEM帶寬限制。所有數(shù)據(jù)類型均存在這一規(guī)律——雙操作數(shù)均在SMEM中的MMA指令，在N<128時(shí)均受SMEM帶寬約束。

2SM MMA實(shí)現(xiàn)了完美的弱擴(kuò)展，相對于1SM MMA在使用兩倍計(jì)算資源時(shí)獲得2倍加速。在SS模式的小形狀配置下，由于操作數(shù)B在兩個(gè)SM間分片，甚至出現(xiàn)超過2倍的加速。研究結(jié)論明確：應(yīng)始終使用給定SMEM tile尺寸下可用的最大指令形狀，以獲得最高吞吐量

延遲方面，所有配置下延遲均隨N從64增至128線性增長，N=256時(shí)出現(xiàn)跳躍。數(shù)據(jù)類型延遲排序呈現(xiàn)規(guī)律性：S8 < BF16 = E4M3 = F4 < MXF8 = MXF4，研究團(tuán)隊(duì)認(rèn)為整數(shù)運(yùn)算功耗效率更高導(dǎo)致S8最快，而微縮放數(shù)據(jù)類型的縮放因子計(jì)算引入了輕微額外開銷。

實(shí)際在途指令數(shù)測試顯示，在典型內(nèi)核使用的1至4條在途MMA指令場景下，4條在途MMA的吞吐量上限約為理論峰值的78%至80%，且1SM MMA比2SM MMA高出約5個(gè)百分點(diǎn)。