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 個百分點。