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

繼《英偉達(dá)張量核心演進(jìn)：從伏特到布萊克威爾》深度文章之后，半導(dǎo)體分析機(jī)構(gòu)SemiAnalysis 投入了數(shù)月工程時(shí)間，深入剖析布萊克威爾架構(gòu)并實(shí)測(cè)原始 PTX 指令性能，以此確立嚴(yán)謹(jǐn)?shù)膶?shí)際性能上限，并與理論峰值進(jìn)行對(duì)比。我們旨在揭示計(jì)算單元與指令級(jí)的硬件吞吐和延遲極限，從機(jī)器學(xué)習(xí)系統(tǒng)與內(nèi)核開發(fā)的角度提供一份實(shí)用的性能刻畫。測(cè)試重點(diǎn)圍繞深度學(xué)習(xí)負(fù)載配置展開，例如對(duì)主流深度學(xué)習(xí)庫 FlashInfer 中采用的異步內(nèi)存拷貝方案進(jìn)行基準(zhǔn)測(cè)試。

布萊克威爾架構(gòu)特性

從霍珀（Hopper）到布萊克威爾（Blackwell），英偉達(dá)對(duì)架構(gòu)進(jìn)行了多項(xiàng)增量改進(jìn)，并針對(duì)與 MMA 相關(guān)的指令調(diào)整了 PTX 抽象層。我們?cè)凇队ミ_(dá)張量核心演進(jìn)》一文中已介紹過其中大部分內(nèi)容。以下是主要的顯著變更：

• 1 引入張量?jī)?nèi)存（TMEM） 用于存儲(chǔ) MMA 累加器。線程不再隱式持有 MMA 運(yùn)算結(jié)果，轉(zhuǎn)而由軟件在 MMA 作用域內(nèi)對(duì) TMEM 進(jìn)行顯式管理。

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

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

集群、GPC 與晶圓良率布局

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

至關(guān)重要的一點(diǎn)：同一個(gè)集群內(nèi)的所有 CTA，保證會(huì)在同一個(gè) GPC（圖形處理集群）上協(xié)同調(diào)度。這一點(diǎn)對(duì) Blackwell 采用 “單 SM 綁定單個(gè) CTA” 的持久 CTA 內(nèi)核模式至關(guān)重要：如果集群大小無法整除 GPC 內(nèi)的 SM 數(shù)量，就會(huì)導(dǎo)致部分 SM 閑置。

這一機(jī)制很容易讓內(nèi)核開發(fā)者困惑：如果不了解文檔記載較少的 GPC 機(jī)制，開發(fā)者往往會(huì)簡(jiǎn)單地按 GPU 總 SM 數(shù)量啟動(dòng)持久 CTA 并開啟集群功能，最終導(dǎo)致部分 CTA 只能串行執(zhí)行。

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

原因是半導(dǎo)體制造過程中會(huì)產(chǎn)生缺陷，這些缺陷可能隨機(jī)出現(xiàn)在芯片的任何位置。因此，英偉達(dá)必須通過架構(gòu)設(shè)計(jì)，讓這些存在缺陷但仍可工作的單元，以相對(duì)統(tǒng)一的方式暴露給軟件使用。

我們通過啟動(dòng)不同大小的集群，并利用 PTX 指令中的%%smid記錄哪些 SM 被分配到同一個(gè) GPC，以此逆向推導(dǎo)出 SM 到 GPC 的映射關(guān)系。

最終得到了 TPC 到 GPC 的邏輯分組列表。這個(gè)列表的長(zhǎng)度超過了 Hopper/Blackwell 標(biāo)配的 8 個(gè) GPC，原因是部分 TPC 會(huì)獨(dú)占一個(gè)邏輯 GPC，永遠(yuǎn)不會(huì)與其他 TPC 協(xié)同調(diào)度。

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

參考資料：

• 集群 API
• 協(xié)作組 API

• CU_LAUNCH_ATTRIBUTE_PREFERRED_CLUSTER_DIMENSION

• CUTLASS 

  示例 73

邏輯 GPC 與物理 GPC

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

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

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

關(guān)鍵術(shù)語注釋

1. quantization issue

   （量化分配問題）

集群大小無法整除 GPC 內(nèi) SM 數(shù)量，導(dǎo)致 SM 閑置的問題

1. preferred cluster size / fallback cluster size

優(yōu)選集群尺寸（高性能優(yōu)先）/ 降級(jí)集群尺寸（兼容性 / 滿資源利用優(yōu)先）

1. logical GPC / physical GPC

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

1. pointer-chase array

指針追蹤數(shù)組，用于精準(zhǔn)測(cè)量緩存訪問延遲的經(jīng)典測(cè)試方法

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

我們同時(shí)用上一節(jié)得出的邏輯 GPC 分組對(duì) SM 進(jìn)行了標(biāo)注；有趣的是，獨(dú)立獨(dú)占的 TPC在延遲上彼此非常接近，且在本次測(cè)試中與 GPC0 高度關(guān)聯(lián)，因此可以推測(cè)這些 TPC 在物理上就位于 GPC0 內(nèi)部。

基于這些數(shù)據(jù)，我們可以進(jìn)一步修正每個(gè) GPC 實(shí)際可用的 TPC 數(shù)量列表，不過其中 5+3 的劃分仍屬于推測(cè)。

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

此外，盡管測(cè)試方式較為間接，我們?nèi)钥傻贸鼋Y(jié)論：裸片間訪問的延遲開銷大約為 300個(gè)時(shí)鐘周期。

這一點(diǎn)在單個(gè) SM 的延遲曲線中也同樣明顯（曲線中同時(shí)包含了大量 L2 擁塞帶來的影響）。

在此特別感謝 Decart AI 的Orian 為本次基準(zhǔn)測(cè)試提供思路啟發(fā)。

存儲(chǔ)子系統(tǒng)

本節(jié)我們介紹存儲(chǔ)子系統(tǒng)，也就是在各個(gè)計(jì)算單元之間搬運(yùn)數(shù)據(jù)的硬件單元。內(nèi)存拷貝指令是使用存儲(chǔ)子系統(tǒng)的核心操作，而新一代架構(gòu)中引入了異步拷貝指令（關(guān)于異步機(jī)制的演進(jìn)可參閱前文）。我們重點(diǎn)關(guān)注兩類異步拷貝指令：LDGSTS和 TMA（張量?jī)?nèi)存加速器）。

異步拷貝

異步拷貝（PTX：cp.async，SASS：LDGSTS）從安培（Ampere）架構(gòu)開始引入，該指令可將數(shù)據(jù)從全局內(nèi)存異步搬運(yùn)到共享內(nèi)存。

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

參考 FlashInfer 的多頭注意力（MHA）內(nèi)核，我們采用以下配置對(duì)異步拷貝進(jìn)行基準(zhǔn)測(cè)試：

• 每個(gè) SM 的 CTA 數(shù)量：1、2、3、4
• 流水線級(jí)數(shù)：1、2、4
• 每個(gè) CTA 的線程數(shù)：64、128、256
• 加載粒度：4B、8B、16B

我們繪制了吞吐率與每個(gè) SM 的飛行字節(jié)數(shù)（即并發(fā)內(nèi)存加載指令正在傳輸?shù)目傋止?jié)數(shù)）的關(guān)系曲線。

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

在相同飛行字節(jié)數(shù)下，16 字節(jié)加載能實(shí)現(xiàn)略高的吞吐，同時(shí)占用更少執(zhí)行資源。例如，在 32 KiB 飛行字節(jié)時(shí)，8B 加載需要 4 級(jí)流水線，而 16B 加載僅需 2級(jí)。這可以節(jié)省兩個(gè)內(nèi)存屏障對(duì)象所需的存儲(chǔ)空間，并降低指令發(fā)射壓力。

整體來看，LDGSTS在 32 KiB 飛行字節(jié)數(shù)下即可達(dá)到飽和，內(nèi)存吞吐約 6.6 TB/s。

我們還針對(duì)實(shí)際 MLA（多層潛在注意力）內(nèi)核常用的配置做了基準(zhǔn)測(cè)試：

• 每個(gè) SM 1 個(gè) CTA

• 16 

  字節(jié)加載
• 每個(gè) CTA 線程數(shù)：64、128、256
• 流水線級(jí)數(shù)：4、8、12、16

實(shí)驗(yàn)表明：增加流水線級(jí)數(shù)能在更高飛行字節(jié)數(shù)下獲得更高吞吐；而提高單個(gè) CTA 的線程數(shù)，在所有配置下都能穩(wěn)定提升性能。

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

我們對(duì)同一組配置進(jìn)行了延遲測(cè)試。結(jié)果顯示：

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

原因在于，為了讓 LDGSTS 達(dá)到高飛行字節(jié)數(shù)，需要啟用大量線程，這會(huì)導(dǎo)致大量線程束（warp）因 MIO（內(nèi)存輸入輸出）節(jié)流 而阻塞。

張量?jī)?nèi)存加速器（TMA）

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

本節(jié)我們以 2D 張量版本（cp.async.bulk.tensor.2d）為代表，測(cè)試 TMA 的典型使用場(chǎng)景性能。

參照 FlashInfer 注意力內(nèi)核的設(shè)置，我們對(duì) TMA 進(jìn)行基準(zhǔn)測(cè)試：每個(gè) SM 只分配一個(gè) CTA，每個(gè) CTA 使用 1~4 個(gè)線程束（warp）中的各一個(gè)線程，來發(fā)起不同塊大小的 TMA 指令。下圖展示了每種飛行字節(jié)數(shù)下的最佳吞吐表現(xiàn)。

本次 TMA 測(cè)試配置如下：

• 每個(gè) SM 的 CTA 數(shù)量：1
• 每個(gè) CTA 的線程數(shù)：128（4     個(gè)線程束）

• TMA 

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

TMA達(dá)到峰值吞吐的時(shí)機(jī)，要比 LDGSTS晚得多。

異步拷貝與 TMA 對(duì)比

像 FlashInfer 這樣的深度學(xué)習(xí)內(nèi)核庫會(huì)同時(shí)使用 TMA 和異步拷貝來加載數(shù)據(jù)。

TMA與異步拷貝具有不同的性能特點(diǎn)：

• TMA 

  適合規(guī)則訪問模式下的大塊數(shù)據(jù)加載，但延遲更高；
• 異步拷貝則能處理不規(guī)則內(nèi)存訪問模式，但存在大小限制。

我們會(huì)說明在不同場(chǎng)景下該如何選擇。本節(jié)針對(duì) FlashInfer 在 MHA 和 MLA 內(nèi)核中實(shí)際使用的配置進(jìn)行了基準(zhǔn)測(cè)試。

可以看到：

• 吞吐方面：飛行字節(jié)數(shù)小于 32 KiB 時(shí)，異步拷貝略優(yōu)于 TMA；超過之后 TMA 迎頭趕上，并且能一直擴(kuò)展到 128      KiB。
• 延遲方面：飛行字節(jié)數(shù)小于 12 KiB 時(shí)，異步拷貝延遲略低于 TMA；超過之后 TMA 延遲會(huì)大幅上升。

在實(shí)際應(yīng)用中，Blackwell 的MLA 內(nèi)核使用異步拷貝來動(dòng)態(tài)加載頁數(shù)據(jù)，而 MHA 內(nèi)核則僅使用 TMA。

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

為了弄清這些內(nèi)核的確切實(shí)現(xiàn)機(jī)制，我們呼吁英偉達(dá)開源 FlashInfer 中的 TRT-LLM 內(nèi)核，以惠及整個(gè)社區(qū)。

TMA組播（Multicast）

TMA支持組播模式：?jiǎn)未渭虞d操作可將數(shù)據(jù)拷貝到多個(gè) SM 的共享內(nèi)存中，目標(biāo)范圍由 CTA 掩碼指定。

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

其核心優(yōu)勢(shì)在于：

• 減少 HBM 讀取，降低有效帶寬占用；
• 顯著減少 L2 流量，因?yàn)槎鄠€(gè) CTA 對(duì)共享數(shù)據(jù)的請(qǐng)求會(huì)被合并為一次請(qǐng)求。

根據(jù) NCU 分析，負(fù)責(zé)處理 TMA 組播請(qǐng)求的硬件單元稱為L2 請(qǐng)求合并器（LRC）：

L2請(qǐng)求合并器（LRC）處理到達(dá) L2 的請(qǐng)求，并在轉(zhuǎn)發(fā)至 L2 緩存前嘗試合并讀請(qǐng)求。

該單元同時(shí)處理來自 SM 的可編程組播請(qǐng)求，并支持寫入壓縮。

這意味著即便不顯式啟用組播，硬件也可能自動(dòng)表現(xiàn)出類似組播的行為，類似于缺失狀態(tài)保持寄存器（MSHR）的機(jī)制。

為驗(yàn)證這一點(diǎn)，我們運(yùn)行了同一套 TMA 組播基準(zhǔn)測(cè)試，但改為所有 CTA 對(duì)同一塊數(shù)據(jù)發(fā)起獨(dú)立 TMA 加載，而非由單個(gè) CTA 執(zhí)行組播加載。

我們對(duì)比了三種場(chǎng)景：

1. 每個(gè) SM 加載不同數(shù)據(jù)（基準(zhǔn)場(chǎng)景）

2. TMA 

   顯式組播 —— 每個(gè)集群中一個(gè) CTA 向集群內(nèi)所有 CTA 執(zhí)行組播加載

3. TMA 

   隱式組播 —— 每個(gè)集群內(nèi)所有 CTA 對(duì)同一塊數(shù)據(jù)執(zhí)行普通 TMA 加載

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

對(duì)于已知的流量模式，顯式 TMA 組播指令可以完全消除冗余 L2 流量，實(shí)現(xiàn)理想的 “每字節(jié) SMEM 數(shù)據(jù)對(duì)應(yīng) 1 / 集群大小 的 L2 數(shù)據(jù)流量”。

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

在有效內(nèi)存吞吐方面，隱式組播與顯式組播表現(xiàn)相當(dāng)。但在降低 L2 緩存流量上，當(dāng)飛行字節(jié)數(shù)超過 64 字節(jié)后，隱式組播的效果就會(huì)明顯下降。

分布式共享內(nèi)存（DSMEM）與本地共享內(nèi)存（SMEM）對(duì)比

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

我們測(cè)試了多種訪問 DSMEM 的PTX 指令模式。在編寫代碼時(shí)，DSMEM 與 SMEM 存在一個(gè)關(guān)鍵區(qū)別：

DSMEM的加載操作是以數(shù)據(jù)包形式傳輸?shù)模c全局內(nèi)存加載類似。因此，DSMEM 的最優(yōu)訪問模式并非本地 SMEM 中避免存儲(chǔ)體沖突的交錯(cuò)訪問，而是更像全局內(nèi)存（GMEM）中典型的連續(xù)合并訪問。

此外我們發(fā)現(xiàn)，若要讓本地 SMEM 達(dá)到 128 字節(jié) / 時(shí)鐘的峰值吞吐，必須使用不帶 ::cluster 修飾符的 ld.shared 指令。

我們?cè)诰帉懟鶞?zhǔn)測(cè)試時(shí)就踩過這個(gè)坑：直接用 ld.shared::cluster 訪問本地和遠(yuǎn)程共享內(nèi)存，結(jié)果無法達(dá)到峰值。

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

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

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

第五代張量核心 MMA

MMA指令是執(zhí)行矩陣乘法的核心操作。從 Hopper 到 Blackwell，MMA性能對(duì)矩陣形狀的依賴性變得越來越強(qiáng)。

本節(jié)我們將深入研究這一現(xiàn)象，通過遍歷不同形狀與數(shù)據(jù)類型，量化分析性能差異。

Blackwell新增了 2SM MMA這一全新類型的 MMA 指令（cta_group::2）：一對(duì) CTA 會(huì)跨兩個(gè) SM 協(xié)同執(zhí)行一次 MMA 運(yùn)算。

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

我們將測(cè)試 2SM MMA 是呈現(xiàn)弱擴(kuò)展、強(qiáng)擴(kuò)展，還是兩者兼具。

我們使用以下配置空間對(duì) MMA 性能進(jìn)行基準(zhǔn)測(cè)試：

吞吐性能

英偉達(dá)針對(duì)不同輸入數(shù)據(jù)類型給出了官方標(biāo)稱吞吐指標(biāo)，本節(jié)我們將展示各類（數(shù)據(jù)格式 + CTA 組）的官方指標(biāo)，并與實(shí)際可達(dá)到的最大吞吐進(jìn)行對(duì)比。

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

吞吐性能

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

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

與此同時(shí)，M=256 在所有配置下都能穩(wěn)定保持接近 100% 的峰值吞吐，原因是 M=256 對(duì)應(yīng)每個(gè) SM 分?jǐn)?/span> M=128，可完整利用數(shù)據(jù)通路。

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

MMA的兩種 AB 布局

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

• SS 

  布局：兩個(gè)輸入矩陣都存放在共享內(nèi)存（SMEM）中

• TS 

  布局：矩陣 A 存放在張量?jī)?nèi)存（TMEM）中，矩陣 B 存放在共享內(nèi)存（SMEM）中

我們觀察到，當(dāng) M=128時(shí)：

• TS 

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

• SS 

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

我們可以證實(shí)，SS 模式下當(dāng) N＜ 128 時(shí)，指令本身會(huì)受到 SMEM 帶寬的限制。

舉個(gè)例子，對(duì)于 FP16 精度：

我們知道每個(gè) SM 每周期硬件可執(zhí)行 8192 MMA FLOPs，而 SMEM 帶寬為 128 字節(jié) / 周期（每 SM）。

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

• A 

  矩陣字節(jié)數(shù) = 2 × M × K =      4096 字節(jié)

• B 

  矩陣字節(jié)數(shù) = 2 × N × K =      2048 字節(jié)
• 浮點(diǎn)運(yùn)算量 FLOPs = 2 × M × N × K = 262144

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

計(jì)算周期 = FLOPs / 16384 = 32 周期

我們逐步增大 N 并計(jì)算后發(fā)現(xiàn)：

只有當(dāng) N ≥ 128 時(shí)，指令才真正進(jìn)入計(jì)算瓶頸階段。

簡(jiǎn)單總結(jié)

• N < 128

  ：受限于共享內(nèi)存（SMEM）帶寬，算力跑不滿

• N ≥ 128

  ：受限于計(jì)算單元算力，達(dá)到理論峰值

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

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

可以清晰看到：N ＜ 262 時(shí)處于內(nèi)存受限區(qū)域，曲線斜率約為128 字節(jié) / 周期，正是 SMEM 的帶寬上限。

2SM MMA在所有數(shù)據(jù)格式與矩陣形狀下均實(shí)現(xiàn)了完美的弱擴(kuò)展：相比 1SM MMA，在計(jì)算資源翻倍的情況下，加速比也恰好達(dá)到 2 倍。

而在 SS 布局的小形狀矩陣中，我們甚至觀察到超過 2 倍的加速比。原因依然是：SS 模式下 N＜128 時(shí)指令受 SMEM 帶寬瓶頸限制，而 2SM 版本會(huì)將操作數(shù) B 分?jǐn)偟絻蓚€(gè)SM 上，從而緩解了帶寬壓力。

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

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

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

延遲

我們對(duì)單條 MMA 指令的延遲進(jìn)行了基準(zhǔn)測(cè)試，并在下圖中對(duì)比展示。

在所有配置下可以看到：延遲從 N=64 到 N=128 呈線性上升，而在 N=256 處出現(xiàn)明顯陡增，這很可能是因?yàn)榫仃嚲S度從 128 躍升到 256 所致。

在單個(gè) CTA 組的 MMA 中：

• 1SM MMA 

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

• 2SM MMA 

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

對(duì)比不同數(shù)據(jù)類型可見：

• 1SM MMA 

  下差異很??；

• 2SM MMA 

  下則出現(xiàn)明顯的延遲分化。

我們觀察到一個(gè)細(xì)微但穩(wěn)定的延遲排序規(guī)律：

S8 < BF16 = E4M3 = F4 < MXF8 = MXF4

我們認(rèn)為，整數(shù)運(yùn)算能效更高，使得 S8 速度最快；而 MXF8 和 MXF4 因?yàn)樾枰~外計(jì)算縮放系數(shù)，引入了少量開銷。

不同飛行指令數(shù)下的吞吐性能

在吞吐基準(zhǔn)測(cè)試中，我們?cè)O(shè)置了大量的飛行指令（256～1024 條）以攤薄指令發(fā)射與提交等待的開銷。

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

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

值得注意的是：

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

對(duì)比 1SM 與 2SM MMA：

• 1SM 

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

在相同數(shù)據(jù)格式與相同 CTA 組配置下：

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

最后我們觀察到：

• 當(dāng)飛行 MMA 指令數(shù)為 4 時(shí)，算力利用率基本封頂在 78%～80%。