英偉達數據中心級Blackwell GPU（SM100）迎來了數代以來幅度最大的 GPU 微架構革新之一，然而官方至今并未發布詳細的技術白皮書。截至目前，面向 AI 負載、針對 UMMA、TMA 等 PTX 與 SASS 指令開展的公開布萊克威爾架構微基準測試研究仍屬空白。

繼《英偉達張量核心演進：從伏特到布萊克威爾》深度文章之后，半導體分析機構SemiAnalysis 投入了數月工程時間，深入剖析布萊克威爾架構并實測原始 PTX 指令性能，以此確立嚴謹的實際性能上限，并與理論峰值進行對比。我們旨在揭示計算單元與指令級的硬件吞吐和延遲極限，從機器學習系統與內核開發的角度提供一份實用的性能刻畫。測試重點圍繞深度學習負載配置展開，例如對主流深度學習庫 FlashInfer 中采用的異步內存拷貝方案進行基準測試。

布萊克威爾架構特性

從霍珀（Hopper）到布萊克威爾（Blackwell），英偉達對架構進行了多項增量改進，并針對與 MMA 相關的指令調整了 PTX 抽象層。我們在《英偉達張量核心演進》一文中已介紹過其中大部分內容。以下是主要的顯著變更：

• 1 引入張量內存（TMEM） 用于存儲 MMA 累加器。線程不再隱式持有 MMA 運算結果，轉而由軟件在 MMA 作用域內對 TMEM 進行顯式管理。

• 2  tcgen05操作現在由單個線程代表整個 CTA（線程塊）發起，而非前代架構中以線程束（warp）或線程束組（warpgroup）為單位。這一點在 CuTe 的 MMA 原子操作中體現明顯：布萊克威爾使用 ThrID =     Layout<_1>，而霍珀基于線程束組的 MMA 則使用 ThrID =      Layout<_128>。

• 3 支持TPC 級別的 TMA以及成對協作 CTA 之間的 MMA，在 PTX 中以 cta_group::2、在 SASS 中以 2CTA 形式暴露。組成一個 TPC 的兩個 SM 可基于共享操作數執行 tcgen05.mma，通過降低單個 CTA 對共享內存（SMEM）的帶寬需求，實現更高運算強度的 MMA 指令。后文將證明，這種操作數共享是充分釋放 MMA 吞吐能力的必要條件。
• 4 原生支持帶微縮放（micro-scaling）的子字節精度數據類型。
• 5 集群啟動控制（CLC）：為持久 CTA 內核中的動態任務調度提供硬件支持（將在后續文章中詳解）。
• 6 程序化依賴啟動（PDL） 在霍珀架構中已引入，用于消除連續內核間的啟動與初始化延遲（將在后續文章中詳解）。
• 
  

集群、GPC 與晶圓良率布局

自 Hopper 架構開始，英偉達數據中心 GPU 就支持一項可選特性，它有多個名稱：線程塊集群、CTA 集群、協作網格陣列（CGA），這些名稱均指向同一功能。集群是 CTA（線程塊）的邏輯分組，其形狀與大小可按每個內核靜態或動態指定。編程模型能感知到集群的存在并實現一些實用功能，例如支持向同一集群內的多個 CTA 執行組播加載—— 我們會在本文后續的 TMA 組播章節詳細講解。

至關重要的一點：同一個集群內的所有 CTA，保證會在同一個 GPC（圖形處理集群）上協同調度。這一點對 Blackwell 采用 “單 SM 綁定單個 CTA” 的持久 CTA 內核模式至關重要：如果集群大小無法整除 GPC 內的 SM 數量，就會導致部分 SM 閑置。

這一機制很容易讓內核開發者困惑：如果不了解文檔記載較少的 GPC 機制，開發者往往會簡單地按 GPU 總 SM 數量啟動持久 CTA 并開啟集群功能，最終導致部分 CTA 只能串行執行。

每個 GPC 最終可用的 SM 數量并非固定值；同一塊芯片上不同 GPC 之間可用 SM 數量不同；甚至同一封裝內的不同裸片之間，可用 SM 布局也可能不對稱。

原因是半導體制造過程中會產生缺陷，這些缺陷可能隨機出現在芯片的任何位置。因此，英偉達必須通過架構設計，讓這些存在缺陷但仍可工作的單元，以相對統一的方式暴露給軟件使用。

我們通過啟動不同大小的集群，并利用 PTX 指令中的%%smid記錄哪些 SM 被分配到同一個 GPC，以此逆向推導出 SM 到 GPC 的映射關系。

最終得到了 TPC 到 GPC 的邏輯分組列表。這個列表的長度超過了 Hopper/Blackwell 標配的 8 個 GPC，原因是部分 TPC 會獨占一個邏輯 GPC，永遠不會與其他 TPC 協同調度。

從 SM100 架構開始，英偉達針對這種量化分配問題提供了解決方案，使內核既能享受大集群帶來的性能優勢，又能充分利用所有可用的 SM 計算單元。啟動內核時可指定兩種集群尺寸：優選集群尺寸與降級集群尺寸。通常情況下，為了完全利用整個 GPU 資源，降級集群尺寸應設置為 2 或 1。

參考資料：

• 集群 API
• 協作組 API

• CU_LAUNCH_ATTRIBUTE_PREFERRED_CLUSTER_DIMENSION

• CUTLASS 

  示例 73

邏輯 GPC 與物理 GPC

我們上文展示的 TPC 到 GPC 的分組屬于邏輯分組。它們僅代表軟件視角下的 GPC 結構，不包含每個 GPC 內部20 個實際物理 SM 中哪些處于啟用狀態的信息，也不體現每個物理 GPC 在雙裸片上的具體位置。

事實上，即便邏輯配置完全相同的 B200 芯片，每個 GPC 中最終可用的物理 SM 數量也不一定完全一致。這可能導致在軟件視角看起來完全相同的 GPU 之間，出現性能不確定性的問題。此外，SM 到 GPC 的邏輯分組信息，也無法區分 B200 封裝內的兩個裸片分別搭載了哪些 GPC。

為了探明 SM 物理布局的更多細節，我們讓每個 SM 遍歷一個指針追蹤數組以填充 L2 緩存，并測量每次加載操作的延遲。針對每個內存地址，我們對比不同 SM 觀測到的加載延遲，最終生成SM 與 SM 之間的距離矩陣（X 軸和Y 軸均為 SM ID）。

關鍵術語注釋

1. quantization issue

   （量化分配問題）

集群大小無法整除 GPC 內 SM 數量，導致 SM 閑置的問題

1. preferred cluster size / fallback cluster size

優選集群尺寸（高性能優先）/ 降級集群尺寸（兼容性 / 滿資源利用優先）

1. logical GPC / physical GPC

邏輯 GPC（軟件看到的分組）/ 物理 GPC（芯片實際硬件布局）

1. pointer-chase array

指針追蹤數組，用于精準測量緩存訪問延遲的經典測試方法

我們可以清晰地看到兩組獨立的 SM 集群，它們之間的 L2 平均訪問延遲相差超過 300 個時鐘周期—— 這顯然就是裸片間（die-to-die）通信的分界。

我們同時用上一節得出的邏輯 GPC 分組對 SM 進行了標注；有趣的是，獨立獨占的 TPC在延遲上彼此非常接近，且在本次測試中與 GPC0 高度關聯，因此可以推測這些 TPC 在物理上就位于 GPC0 內部。

基于這些數據，我們可以進一步修正每個 GPC 實際可用的 TPC 數量列表，不過其中 5+3 的劃分仍屬于推測。

• 裸片 A：[10, 10, 10, 9]
• 裸片 B：[9, 9, 9, 5+3]

此外，盡管測試方式較為間接，我們仍可得出結論：裸片間訪問的延遲開銷大約為 300個時鐘周期。

這一點在單個 SM 的延遲曲線中也同樣明顯（曲線中同時包含了大量 L2 擁塞帶來的影響）。

在此特別感謝 Decart AI 的Orian 為本次基準測試提供思路啟發。

存儲子系統

本節我們介紹存儲子系統，也就是在各個計算單元之間搬運數據的硬件單元。內存拷貝指令是使用存儲子系統的核心操作，而新一代架構中引入了異步拷貝指令（關于異步機制的演進可參閱前文）。我們重點關注兩類異步拷貝指令：LDGSTS和 TMA（張量內存加速器）。

異步拷貝

異步拷貝（PTX：cp.async，SASS：LDGSTS）從安培（Ampere）架構開始引入，該指令可將數據從全局內存異步搬運到共享內存。

異步拷貝是非阻塞的，允許內存加載與計算操作并行執行。它還能直接寫入共享內存，無需經過寄存器，從而降低寄存器占用壓力。

參考 FlashInfer 的多頭注意力（MHA）內核，我們采用以下配置對異步拷貝進行基準測試：

• 每個 SM 的 CTA 數量：1、2、3、4
• 流水線級數：1、2、4
• 每個 CTA 的線程數：64、128、256
• 加載粒度：4B、8B、16B

我們繪制了吞吐率與每個 SM 的飛行字節數（即并發內存加載指令正在傳輸的總字節數）的關系曲線。

盡管不同加載粒度在相同飛行字節數下最終能達到相近的吞吐率，但我們更推薦使用 16字節加載。

在相同飛行字節數下，16 字節加載能實現略高的吞吐，同時占用更少執行資源。例如，在 32 KiB 飛行字節時，8B 加載需要 4 級流水線，而 16B 加載僅需 2級。這可以節省兩個內存屏障對象所需的存儲空間，并降低指令發射壓力。

整體來看，LDGSTS在 32 KiB 飛行字節數下即可達到飽和，內存吞吐約 6.6 TB/s。

我們還針對實際 MLA（多層潛在注意力）內核常用的配置做了基準測試：

• 每個 SM 1 個 CTA

• 16 

  字節加載
• 每個 CTA 線程數：64、128、256
• 流水線級數：4、8、12、16

實驗表明：增加流水線級數能在更高飛行字節數下獲得更高吞吐；而提高單個 CTA 的線程數，在所有配置下都能穩定提升性能。

有意思的是，MLA 內核采用 2 個線程束（warp）+ 12 級流水線，實測吞吐約 2.2 TB/s。我們認為原因在于：執行 softmax 的線程束需要占用大量寄存器，增加線程束數量會導致單個線程可分配的寄存器減少，從而限制性能。

我們對同一組配置進行了延遲測試。結果顯示：

LDGSTS的基線延遲約為600 納秒，并且在飛行字節數超過 8 KiB 后，延遲幾乎翻倍。

原因在于，為了讓 LDGSTS 達到高飛行字節數，需要啟用大量線程，這會導致大量線程束（warp）因 MIO（內存輸入輸出）節流 而阻塞。

張量內存加速器（TMA）

TMA（PTX 指令：cp.async.bulk.tensor，SASS 指令：UTMALDG）是在 Hopper 架構中引入的異步數據拷貝引擎，專門用于將大量數據從全局內存搬運到共享內存。只需單個線程即可發起 TMA 操作，完成地址生成、內存交織（swizzling）與越界處理，從而讓其他線程可以執行獨立任務。

本節我們以 2D 張量版本（cp.async.bulk.tensor.2d）為代表，測試 TMA 的典型使用場景性能。

參照 FlashInfer 注意力內核的設置，我們對 TMA 進行基準測試：每個 SM 只分配一個 CTA，每個 CTA 使用 1~4 個線程束（warp）中的各一個線程，來發起不同塊大小的 TMA 指令。下圖展示了每種飛行字節數下的最佳吞吐表現。

本次 TMA 測試配置如下：

• 每個 SM 的 CTA 數量：1
• 每個 CTA 的線程數：128（4     個線程束）

• TMA 

  塊維度：2D 尺寸從     32×8 逐步增大到 128×128

TMA達到峰值吞吐的時機，要比 LDGSTS晚得多。

異步拷貝與 TMA 對比

像 FlashInfer 這樣的深度學習內核庫會同時使用 TMA 和異步拷貝來加載數據。

TMA與異步拷貝具有不同的性能特點：

• TMA 

  適合規則訪問模式下的大塊數據加載，但延遲更高；
• 異步拷貝則能處理不規則內存訪問模式，但存在大小限制。

我們會說明在不同場景下該如何選擇。本節針對 FlashInfer 在 MHA 和 MLA 內核中實際使用的配置進行了基準測試。

可以看到：

• 吞吐方面：飛行字節數小于 32 KiB 時，異步拷貝略優于 TMA；超過之后 TMA 迎頭趕上，并且能一直擴展到 128      KiB。
• 延遲方面：飛行字節數小于 12 KiB 時，異步拷貝延遲略低于 TMA；超過之后 TMA 延遲會大幅上升。

在實際應用中，Blackwell 的MLA 內核使用異步拷貝來動態加載頁數據，而 MHA 內核則僅使用 TMA。

FlashInfer中大部分Blackwell MHA 內核均由 TRT-LLM 貢獻，因此我們只能通過反匯編二進制文件來推測內核邏輯。我們發現，與 Hopper 類似，所有 Blackwell TRT-LLM 內核都使用 TMA。我們推測，在動態頁加載場景下，這些內核沿用了 Hopper 的設計思路：使用4D TMA，將頁索引作為最后一維，并在需要時通過 TensorMap 進行索引尋址。

為了弄清這些內核的確切實現機制，我們呼吁英偉達開源 FlashInfer 中的 TRT-LLM 內核，以惠及整個社區。

TMA組播（Multicast）

TMA支持組播模式：單次加載操作可將數據拷貝到多個 SM 的共享內存中，目標范圍由 CTA 掩碼指定。

組播常用于 GEMM 類計算模式 ——多個 SM 處理不同輸出分塊時，輸入分塊可在 SM 間共享。例如在 SwiGLU 激活函數中，兩個 GEMM 操作共用同一個輸入矩陣，組播就非常適用。

其核心優勢在于：

• 減少 HBM 讀取，降低有效帶寬占用；
• 顯著減少 L2 流量，因為多個 CTA 對共享數據的請求會被合并為一次請求。

根據 NCU 分析，負責處理 TMA 組播請求的硬件單元稱為L2 請求合并器（LRC）：

L2請求合并器（LRC）處理到達 L2 的請求，并在轉發至 L2 緩存前嘗試合并讀請求。

該單元同時處理來自 SM 的可編程組播請求，并支持寫入壓縮。

這意味著即便不顯式啟用組播，硬件也可能自動表現出類似組播的行為，類似于缺失狀態保持寄存器（MSHR）的機制。

為驗證這一點，我們運行了同一套 TMA 組播基準測試，但改為所有 CTA 對同一塊數據發起獨立 TMA 加載，而非由單個 CTA 執行組播加載。

我們對比了三種場景：

1. 每個 SM 加載不同數據（基準場景）

2. TMA 

   顯式組播 —— 每個集群中一個 CTA 向集群內所有 CTA 執行組播加載

3. TMA 

   隱式組播 —— 每個集群內所有 CTA 對同一塊數據執行普通 TMA 加載

TMA組播能夠提供極高的加載帶寬以填充 SMEM 緩沖區，即便數據尚未緩存到 L2 中也是如此。

對于已知的流量模式，顯式 TMA 組播指令可以完全消除冗余 L2 流量，實現理想的 “每字節 SMEM 數據對應 1 / 集群大小 的 L2 數據流量”。

我們還觀察到，在這一簡單測試中，顯式與隱式模式下 SMEM 填充帶寬幾乎一致。但 LRC 并非完美：隱式模式下 L2 仍會產生略多的流量，尤其在總數據量增大時更為明顯。

在有效內存吞吐方面，隱式組播與顯式組播表現相當。但在降低 L2 緩存流量上，當飛行字節數超過 64 字節后，隱式組播的效果就會明顯下降。

分布式共享內存（DSMEM）與本地共享內存（SMEM）對比

英偉達在 Hopper 架構中引入了分布式共享內存（DSMEM）。DSMEM 允許同一個集群內的線程塊（CTA）互相訪問彼此的共享內存，這對CTA 間歸約等計算模式非常實用。但通過 DSMEM 讀取對等 CTA 內存的吞吐率，遠低于本地 SMEM 每時鐘周期 128 字節的峰值。

我們測試了多種訪問 DSMEM 的PTX 指令模式。在編寫代碼時，DSMEM 與 SMEM 存在一個關鍵區別：

DSMEM的加載操作是以數據包形式傳輸的，與全局內存加載類似。因此，DSMEM 的最優訪問模式并非本地 SMEM 中避免存儲體沖突的交錯訪問，而是更像全局內存（GMEM）中典型的連續合并訪問。

此外我們發現，若要讓本地 SMEM 達到 128 字節 / 時鐘的峰值吞吐，必須使用不帶 ::cluster 修飾符的 ld.shared 指令。

我們在編寫基準測試時就踩過這個坑：直接用 ld.shared::cluster 訪問本地和遠程共享內存，結果無法達到峰值。

• 使用     ld.shared 時，編譯器會生成專用的 LDS 指令；
• 而使用     ld.shared::cluster 時，編譯器只會生成通用的 LD     指令，無法讓本地 SMEM 跑出峰值性能。

我們還發現，ld.shared::cluster 很難進一步提升吞吐；只有切換為cp.async.bulk（PTX）/ UBLKCP（SASS） 后，才能通過單指令搬運更大數據量，讓 DSMEM 吞吐獲得小幅提升。

以下是我們使用不同 PTX 模式測得的峰值吞吐，單位為字節 / 時鐘周期（B/clk），便于與本地 SMEM 的理論最大值對比。

第五代張量核心 MMA

MMA指令是執行矩陣乘法的核心操作。從 Hopper 到 Blackwell，MMA性能對矩陣形狀的依賴性變得越來越強。

本節我們將深入研究這一現象，通過遍歷不同形狀與數據類型，量化分析性能差異。

Blackwell新增了 2SM MMA這一全新類型的 MMA 指令（cta_group::2）：一對 CTA 會跨兩個 SM 協同執行一次 MMA 運算。

具體來說，輸入矩陣 A 會被復制，矩陣B 和矩陣 D 則在兩個 SM 間分片，并且這對 CTA 可以互相訪問彼此的共享內存。這使得更大規模的 MMA 形狀成為可能。

我們將測試 2SM MMA 是呈現弱擴展、強擴展，還是兩者兼具。

我們使用以下配置空間對 MMA 性能進行基準測試：

吞吐性能

英偉達針對不同輸入數據類型給出了官方標稱吞吐指標，本節我們將展示各類（數據格式 + CTA 組）的官方指標，并與實際可達到的最大吞吐進行對比。

結果表明，UMMA 在所有格式與CTA 組配置下均能實現接近理論峰值的吞吐，即便在需要協同開銷的 2SM 版本上也是如此。

吞吐性能

在所有 N 尺寸下的 1SM MMA 配置中可以看到：較小的 M=64 僅能達到理論峰值吞吐的 50%，而更大的 M=128 則能接近 100%。這證實 M=64 只利用了一半的數據通路。

在 2SM MMA 上，M=128 在 N=64 時吞吐約為峰值的 90%，在其余所有 N 尺寸下均接近 100%。M128N64的吞吐瓶頸應來自 TMEM、L2、SMEM 等其他硬件單元。

與此同時，M=256 在所有配置下都能穩定保持接近 100% 的峰值吞吐，原因是 M=256 對應每個 SM 分攤 M=128，可完整利用數據通路。

我們還發現：相同位寬的數據格式，吞吐表現完全一致；采用微縮放的數據類型幾乎沒有額外開銷。

MMA的兩種 AB 布局

MMA支持兩種不同的 AB 矩陣存儲布局：

• SS 

  布局：兩個輸入矩陣都存放在共享內存（SMEM）中

• TS 

  布局：矩陣 A 存放在張量內存（TMEM）中，矩陣 B 存放在共享內存（SMEM）中

我們觀察到，當 M=128時：

• TS 

  布局在所有 N 尺寸下都能達到接近峰值的吞吐；

• SS 

  布局在 N 較小時性能偏低，直到 N=128時才追平峰值性能。

我們可以證實，SS 模式下當 N＜ 128 時，指令本身會受到 SMEM 帶寬的限制。

舉個例子，對于 FP16 精度：

我們知道每個 SM 每周期硬件可執行 8192 MMA FLOPs，而 SMEM 帶寬為 128 字節 / 周期（每 SM）。

以 M=128、N=64、K=16為例：

• A 

  矩陣字節數 = 2 × M × K =      4096 字節

• B 

  矩陣字節數 = 2 × N × K =      2048 字節
• 浮點運算量 FLOPs = 2 × M × N × K = 262144

SMEM訪問周期 = (A_bytes + B_bytes) / 128 = 48 周期

計算周期 = FLOPs / 16384 = 32 周期

我們逐步增大 N 并計算后發現：

只有當 N ≥ 128 時，指令才真正進入計算瓶頸階段。

簡單總結

• N < 128

  ：受限于共享內存（SMEM）帶寬，算力跑不滿

• N ≥ 128

  ：受限于計算單元算力，達到理論峰值

其他數據類型也是同理 ——兩個操作數都放在 SMEM 中的 MMA 指令，在 N 小于 128 時均受 SMEM 帶寬瓶頸限制。

為進一步說明這一點，我們繪制了 FP8 精度下 1SM MMA 所有形狀的屋頂線曲線。

可以清晰看到：N ＜ 262 時處于內存受限區域，曲線斜率約為128 字節 / 周期，正是 SMEM 的帶寬上限。

2SM MMA在所有數據格式與矩陣形狀下均實現了完美的弱擴展：相比 1SM MMA，在計算資源翻倍的情況下，加速比也恰好達到 2 倍。

而在 SS 布局的小形狀矩陣中，我們甚至觀察到超過 2 倍的加速比。原因依然是：SS 模式下 N＜128 時指令受 SMEM 帶寬瓶頸限制，而 2SM 版本會將操作數 B 分攤到兩個SM 上，從而緩解了帶寬壓力。

SS模式：當 N ＜ 128 時，由于受 SMEM 帶寬瓶頸限制，加速比超過 2 倍。

TS模式：接近完美的2 倍加速。

這些實驗表明：在給定的 SMEM 分塊大小下，想要獲得最大吞吐，應始終使用盡可能大的指令形狀。

延遲

我們對單條 MMA 指令的延遲進行了基準測試，并在下圖中對比展示。

在所有配置下可以看到：延遲從 N=64 到 N=128 呈線性上升，而在 N=256 處出現明顯陡增，這很可能是因為矩陣維度從 128 躍升到 256 所致。

在單個 CTA 組的 MMA 中：

• 1SM MMA 

  的 M=64 和 M=128 在各 N 尺寸下延遲相近；

• 2SM MMA 

  中，M=256 的延遲增長略快于 M=128，這與我們的理論估算一致。

對比不同數據類型可見：

• 1SM MMA 

  下差異很小；

• 2SM MMA 

  下則出現明顯的延遲分化。

我們觀察到一個細微但穩定的延遲排序規律：

S8 < BF16 = E4M3 = F4 < MXF8 = MXF4

我們認為，整數運算能效更高，使得 S8 速度最快；而 MXF8 和 MXF4 因為需要額外計算縮放系數，引入了少量開銷。

不同飛行指令數下的吞吐性能

在吞吐基準測試中，我們設置了大量的飛行指令（256～1024 條）以攤薄指令發射與提交等待的開銷。

但實際內核通常只使用 1～4 條飛行中的 MMA 指令。因此我們專門測試了飛行指令數為 1～10 時的吞吐，并分析其變化規律。

在所有配置下，相同 N 值與相同MMA 飛行數所達到的理論算力利用率（SoL, Speed-of-Light） 比例相近。

值得注意的是：

• 只有最大的 N 尺寸能達到 90% 左右的算力利用率；
• 最小的 N 尺寸僅能達到約 70%。

對比 1SM 與 2SM MMA：

• 1SM 

  的算力利用率比 2SM 高出約 5%。

在相同數據格式與相同 CTA 組配置下：

• 更大的 N 尺寸，吞吐始終高于更小的 N。

最后我們觀察到：

• 當飛行 MMA 指令數為 4 時，算力利用率基本封頂在 78%～80%。