公眾號記得加星標⭐️，第一時間看推送不會錯過。
來源：內容編譯自semianalysis。
在人工智慧和深度學習領域，GPU 計算能力的提升速度遠超摩爾定律，年復一年地持續實現著“黃氏定律”般顯著的效能提升。推動這一進步的核心技術正是 Tensor Core。
儘管 Tensor Core 無疑是現代人工智慧和機器學習的基石，但即使是許多經驗豐富的從業者，對其也仍未有深入的理解。GPU 架構以及基於該架構的程式設計模型的快速發展，使得機器學習研究人員和科學家越來越難以跟上 Tensor Core 的最新變化並理解這些變化的影響。
在本問中，我們將介紹主流資料中心 GPU 的核心特性，首先解釋效能工程的重要基本原理。然後，我們將追溯 Nvidia Tensor Core 架構和程式設計模型的演變，並重點闡述其演變背後的動機。我們的最終目標是提供資源，幫助理解 Nvidia 的 GPU 架構，並直觀地瞭解其架構的演變。只有在解釋完每個架構之後，我們才能解釋 Blackwell 張量核心及其全新記憶體層次結構的精妙之處。
需要強調的是，紮實的計算機架構理解能力是理解本文諸多講解和討論的先決條件。本文將簡要介紹 CUDA 程式設計，供讀者複習，而非解釋 GPU 架構的基本概念。相反，我們將以 Tensor Core 的前沿知識為基礎，透過詳盡的講解，將目前較為零散的知識轉化為易於理解的結構化見解，從而拓展讀者對這項前沿技術的理解。
效能優先原則
阿姆達爾定律（Amdahl’s Law）
對於固定的問題規模，阿姆達爾定律規定了透過增加計算資源進行並行化可以獲得的最大加速比。具體而言，擴充套件計算資源只會縮短並行部分的執行時間，因此效能提升受限於序列部分的執行時間。為了量化這一點，最大效能提升如下：
其中 S 是並行工作的執行時間，p 是可並行化工作的加速比。在並行部分完全並行化的理想情況下，加速比 p 可以是處理單元的數量。
強擴充套件和弱擴充套件
（Strong and Weak Scaling）
強擴充套件和弱擴充套件描述了針對不同問題設定擴充套件計算資源的效能提升。強擴充套件是指擴充套件計算資源以解決固定規模的問題，而阿姆達爾定律則量化了強擴充套件的加速比。另一方面，弱擴充套件是指擴充套件計算資源以在恆定時間內解決更大的問題。例如，使用 4 倍的計算資源，在相同的時間內處理 4 倍大的影像。
強擴充套件和弱擴充套件在不同規模的問題上意味著不同的效能提升。強擴充套件可以為所有規模的問題提供加速，而弱擴充套件僅在使用更多計算來解決更大問題時才保證效能提升。
資料移動是大忌
資料移動是一種罪過，因為就執行時間和擴充套件性而言，計算成本低廉，而資料移動成本高昂。資料移動速度從根本上來說更慢，因為現代 DRAM 單元的執行速度為數十納秒，而電晶體的開關速度為亞納秒。就擴充套件性而言，雖然自 2000 年代以來計算速度的提升有所放緩，但記憶體速度的提升也更慢，從而形成了“記憶體牆”。
Tensor Core 架構演進
張量核心生成概述
在本節中，我們將介紹使用 Tensor Core 的主要 Nvidia GPU 架構，即 Tesla V100 GPU、A100 Tensor Core GPU、H100 Tensor Core GPU 以及 Blackwell GPU。我們還添加了一個 Tensor Core 之前的章節，作為 CUDA 程式設計模型的複習。我們將簡要介紹與理解 Tensor Core 相關的主要特性和變化，並將詳細資訊保留到其他來源，我們會在每個小節中提供連結。
預張量核心
一、PTX程式設計模型
並行執行緒執行 (PTX：Parallel Thread Execution) 是跨 GPU 代抽象的虛擬指令集。PTX 程式描述了一個核函式，該函式由大量 GPU 執行緒執行，這些執行緒在 GPU 的硬體執行單元（即 CUDA 核心）上執行。執行緒被組織為網格，每個網格由協作執行緒陣列 ( CTA )組成。PTX 執行緒可以訪問來自多個狀態空間的資料，這些狀態空間是具有不同特性的記憶體儲存區域。具體而言，執行緒具有每個執行緒的暫存器，CTA 內的執行緒具有共享記憶體，並且所有執行緒都可以訪問全域性記憶體。
二、PTX 機器模型
GPU 架構圍繞流多處理器 ( SM )陣列構建。SM 由標量處理核心、多執行緒指令單元和片上共享記憶體組成。SM 將每個執行緒對映到一個標量處理核心（也稱為 CUDA 核心），而多執行緒指令單元則以 32 個並行執行緒組（稱為Warp）的形式管理執行緒。
在指令發出時，指令單元選擇一個 Warp，並向 Warp 中的執行緒發出指令。這種執行方法稱為單指令多執行緒 ( SIMT )。與單指令多資料 ( SIMD ) 類似，SIMT 使用一條指令控制多個處理單元，但與 SIMD 不同的是，SIMT 指定的是單執行緒行為，而不是向量寬度。
三、流式彙編器
流式彙編器 (SASS：Streaming Assembler ) 是 PTX 虛擬化所基於的特定於架構的指令集。有關更多資訊，請參閱CUDA 二進位制實用程式文件。遺憾的是，由於 NVIDIA 向競爭對手隱藏了其架構 ISA 的細節，SASS 的文件並不完善。
Volta
NVIDIA 為何新增 Tensor Core
隨著深度學習變得越來越突出，業界注意到 ML 工作負載需要硬體加速。2015 年初，Google 部署了 TPUv1 來加速其內部 ML 工作負載，2017 年，Nvidia 推出了用於矩陣數學的專用硬體。雖然 GPU 由於其簡單的硬體流水線在發出指令時會消耗少量能量（~30pJ），但簡單的浮點運算消耗的HFMA 能量甚至更少，僅為 1.5pJ。這使得指令所需的功耗是浮點運算本身的 20 倍。因此，執行大量浮點運算進行矩陣乘法是功耗低的。為了攤銷指令開銷，我們需要使用每個指令可以執行更多計算的複雜指令。為此，Nvidia 設計了半精度矩陣乘法和累加（HMMA）指令，這是一條執行半精度矩陣乘法的專用指令。執行該指令的相應專用硬體是 Tensor Core，它於 2017 年在 Volta 架構的 Tesla V100 GPU 中推出。Volta 張量核心是在 Volta 架構開發的後期新增的，僅在流片前幾個月，這證明了 Nvidia 對其架構的調整速度有多快。
MMA 指導概述
給定一個矩陣，乘法和累加 (MMA) 指令計算 D = A * B + C：
A 是 M×K 矩陣
B 是 K×N 矩陣
C 和 D 是 M×N 矩陣
我們將矩陣形狀表示為mMnNkK或 MxNxK。
為了執行完整的計算，我們首先將矩陣 A、B 和 C 從共享記憶體載入到執行緒暫存器，以便每個執行緒儲存矩陣的片段。其次，我們執行 MMA 指令，該指令從執行緒暫存器讀取矩陣，在 Tensor Core 上執行計算，並將結果儲存到執行緒暫存器。最後，我們將結果從執行緒暫存器儲存回共享記憶體。完整的計算由多個執行緒共同執行，這意味著每個步驟都需要協作執行緒之間的同步。
第一代 Tensor Core – Warp-scoped MMA
Tesla V100 GPU 的 SM 包含 8 個 Tensor Core，分為兩部分。每個 Tensor Core 每週期能夠計算相當於 4x4x4 矩陣乘法的運算，相當於每個 SM 每週期 1024 FLOP。
NVIDIA 設計了 PTX 指令 mma，以針對較低級別的HMMA指令。在 Volta 架構上，MMA 指令執行 8x8x4 矩陣乘法，由 8 個執行緒組成的四對 (quadpair) 透過共同儲存輸入和輸出矩陣參與運算。其中，T0 表示執行緒 0，[T0, T1, T2, T3] 和 [T16, T17, T18, T19] 表示執行緒組，這兩個執行緒組組成一個四對 (quadpair)。
在資料型別方面，Volta Tensor Core 支援 FP16 輸入和 FP32 累積，這與 NVIDIA 的混合精度訓練技術相呼應。該技術表明，可以在不損失模型精度的情況下以較低的精度訓練模型。
Turing
Turing 架構包含第二代 Tensor Core，這是 Volta Tensor Core 的增強版，增加了對 INT8 和 INT4 精度的支援。Turing Tensor Core 支援全新的 Warp-Level 同步 MMA，我們將在下一節中討論。Turing Tensor Core 還支援深度學習超級取樣 (DLSS)，標誌著 NVIDIA 開始將深度學習應用於遊戲圖形。
Ampere
非同步資料複製
NVIDIA 在 Ampere 架構中引入了非同步資料複製，這是一種以非同步方式將資料直接從全域性記憶體複製到共享記憶體的方法。要在 Volta 架構上將資料從全域性記憶體載入到共享記憶體，執行緒必須先將資料從全域性記憶體載入到暫存器，然後再將其儲存到共享記憶體。然而，MMA 指令的暫存器使用率很高，並且必須與資料載入操作共享暫存器檔案，這會導致暫存器壓力過大，並浪費記憶體頻寬來複制資料進出 RF。
非同步資料複製透過從全域性記憶體 (DRAM) 獲取資料並將其直接儲存到共享記憶體（可選 L1 訪問）來緩解此問題，從而釋放更多暫存器用於 MMA 指令。資料載入和計算可以非同步進行，這從程式設計模型的角度來看更加困難，但可以提高效能。
此功能透過 PTX 指令執行緒級非同步複製 cp.async 實現（文件）。對應的 SASS 是 LDGSTS，即非同步全域性到共享記憶體複製。具體的同步方法是非同步組和基於 mbarrier 的完成機制。
第三代 Tensor Core – 曲速級同步
MMA（Warp-level Synchronous MMA）
Ampere 每個 SM 有 4 個 Tensor Core，每個 Tensor Core 每週期能夠執行 512 FLOP，每個 SM 每週期總計 2048 Dense FLOP，效能是 Volta 的兩倍。
Volta 需要 8 個執行緒組成的四對才能參與 MMA 運算，而 Ampere 則需要 32 個執行緒的完整 Warp。採用 Warp 寬度的 MMA 指令簡化了執行緒佈局，並降低了 Ampere 的 RF 壓力。例如，以下是 16x8x16 形狀的混合精度浮點的執行緒和資料佈局：
NVIDIAldmatrix在 Ampere 中引入了增強型向量化載入操作。與類似mma，ldmatrix它也是 Warp 範圍的，這意味著一組 Warp 執行緒共同載入一個矩陣。與發出多個載入指令相比，這減少了地址生成暫存器的使用，從而降低了暫存器壓力。有關更多資訊，請參閱CUDA 文件。
ldmatrix將資料載入到暫存器中，其佈局與 Tensor Core 的資料佈局相匹配。與 Volta 的交錯模式相比（參見“Tensor Core 程式設計：使用 CUTLASS 實現原生 Tensor Core”），更簡單的執行緒和資料佈局極大地提升了程式設計的人體工程學。觀看 GTC 演講“開發 CUDA 核心以在 NVIDIA A100 上將 Tensor Core 推向極限”。
Ampere MMA 採用 Brain 浮點格式 (BF16)，該格式已成為半精度資料型別的事實標準。BF16 提供與 FP32 相同的 8 位指數範圍，但尾數為 7 位，從而能夠以一半的儲存成本實現 FP32 級別的動態範圍。BF16 還消除了混合精度訓練中損失縮放的需要。
Hopper
執行緒塊叢集
隨著 SM 數量的增加，單個 SM 與整個 GPU 之間的大小差異也隨之增大。為了在 CTA（對映到 SM）和網格（對映到整個 GPU）之間提供更精細的控制粒度，NVIDIA 在 Hopper 上添加了一個新的執行緒層次結構級別——執行緒塊叢集，它對映到物理上位於同一圖形處理叢集 (GPC) 中的一組 SM。執行緒塊叢集也稱為協作網格陣列 (CGA)，在 CUDA 文件中簡稱為叢集。
執行緒塊叢集中的 CTA 保證在同一 GPC 內的各個 SM 上協同排程，並且預設每個 SM 分配一個 CTA。這些 SM 的共享記憶體分割槽構成分散式共享記憶體 (DSMEM)。執行緒可以透過專用的 SM 到 SM 網路（無需經過 L2 快取）以低延遲訪問其他 SM 的共享記憶體。透過將 GPC 硬體執行單元暴露給程式設計模型，程式設計師可以減少資料移動並提高資料區域性性。
張量記憶加速器
為了提高資料獲取效率，NVIDIA 為每個 Hopper SM 添加了張量記憶體加速器 (TMA)。TMA 是一個專用硬體單元，可加速全域性記憶體和共享記憶體之間的大量非同步資料傳輸（批次非同步複製）。
CTA 中的單個執行緒可以啟動 TMA 複製操作。TMA 釋放執行緒來執行其他獨立工作，處理地址生成並提供額外的優勢，例如越界處理。在 PTX 中，相應的指令是cp.async.bulk，詳情請參閱CUDA 文件中的 部分。
然而，對於小型請求，由於地址生成開銷，TMA 載入的延遲比常規非同步資料複製更高。因此，NVIDIA 建議程式設計師使用 TMA 進行大型資料複製，以分攤開銷。例如，在 LLM 推理中，TMA 不適用於以小塊載入鍵值快取的工作負載，但當每個塊是 16 位元組的倍數時，TMA 效果良好。
TMA 還支援一種稱為多播的資料載入模式，該模式將資料從全域性記憶體載入到執行緒塊叢集中多個 SM 的共享記憶體中，由多播掩碼指定。多播載入不是發出多個全域性記憶體載入請求，將同一段資料載入到多個 SM 中，而是一次性完成。具體來說，執行緒塊叢集中的多個 CTA 將部分資料載入到其對應的 SMEM 中，並透過 DSMEM 共享資料。這減少了二級快取流量，進而減少了 HBM 流量。
第四代 Tensor Core – Warpgroup
級非同步 MMA
NVIDIA 透過 Hopper 引入了一種新型的 MMA，即 Warpgroup 級 MMA（wgmma）。wgmma它是 Warpgroup 級的，這意味著由 4 個 Warp 組成的 Warpgroup 共同執行 MMA 操作。wgmma支援更廣泛的形狀。例如，混合精度 MMA 支援m64nNk16，其中 N 可以是 8 的倍數，範圍從 8 到 256。wgmma.mma_async降低到一組新的 SASS：GMMA。在另一個示例中，半精度wgmma指令降低到HGMMA。
雖然 Warpgroup 中的所有執行緒都會將輸出矩陣儲存在其暫存器中，但 Hopper Tensor Core 可以直接從共享記憶體（而非暫存器）載入運算元，從而節省暫存器空間和頻寬。具體來說，運算元矩陣 A 可以駐留在暫存器或共享記憶體中，而運算元矩陣 B 只能透過共享記憶體訪問。
在wgmma資料型別方面，Hopper 引入了 8 位浮點資料型別（E4M3 和 E5M2），並採用 FP32 累加。實際上，累加路徑採用 22 位定點格式（13 位尾數加符號位和指數位），與真正的 32 位累加相比，動態範圍有所限制。由於張量核心精度降低，為了避免影響訓練精度，每次 N_c 次累加都必須在 CUDA 核心中進行。。這種降低精度的累加提高了效率，但卻以犧牲精度為代價。
Blackwell
巨大的暫存器壓力並沒有讓 Hopper 放鬆，這促使Tensor Memory (TMEM)應運而生，這是一種專門用於 Tensor Core 運算的新型記憶體。在每個 SM 上，TMEM 擁有 128 行（通道）和 512 列 4 位元組單元，總計 256 KB，這也是 SM 上暫存器檔案的大小。
TMEM 的記憶體訪問模式受到限制。具體來說，訪問整個 TMEM 需要一個 WarpGroup，並且 WarpGroup 中的每個 Warp 只能訪問一組特定的通道。透過限制記憶體訪問模式，硬體設計人員可以減少訪問埠的數量，從而節省晶片空間。另一方面，這種設計也意味著結語操作也需要一個 WarpGroup 來執行。與共享記憶體不同，程式設計師必須明確管理 TMEM，包括分配、釋放以及將資料複製到 TMEM 和從 TMEM 複製資料。
CTA Pair
如果執行緒塊叢集中的兩個 CTA 在其執行緒塊叢集中的 CTA 排序最後一位不同（例如 0 和 1、4 和 5），則它們會形成一個CTA 對。一個 CTA 對對映到一個紋理處理叢集 (TPC)，該叢集由兩個 SM 組成，並與其他 TPC 組合形成一個 GPC。當 Blackwell Tensor Core 運算以 CTA 對粒度執行時，這兩個 CTA 能夠共享輸入運算元。這種共享可以降低 SMEM 的容量和頻寬需求。
Tensor Core 第五代 MMA
Tensor Core 第五代 MMA 指令（tcgen05.mma在 PTX 中）已完全不再使用暫存器來儲存矩陣。運算元現在駐留在共享記憶體和 Tensor Memory 中。
具體來說，假設 MMA 計算 D = A * B + D：不使用執行緒暫存器可以消除複雜的資料佈局，並釋放執行緒暫存器空間用於其他工作，例如尾聲操作。與wgmma使用 Warpgroup 啟動 MMA 操作不同，tcgen05.mma它具有單執行緒語義，這意味著單個執行緒可以啟動 MMA 操作。這消除了 Warp 發出 MMA 操作的作用。
一個值得注意的 MMA 變體是 MMA.2SM，它使用兩個 SM 共同執行 MMA 操作。
MMA.2SM 以 CTA 對級粒度執行，並且由於其tcgen05.mma具有單執行緒語義，因此 CTA 對中領導者 CTA 中的單個執行緒將啟動 MMA.2SM。這裡我們展示了資料路徑組織布局 A。佈局 A 顯示，與 1SM 版本（佈局 D）相比，MMA.2SM 將 M 維度增加了一倍，因此這兩個 SM 載入不同的矩陣 A 和 D 分塊。此外，MMA.2SM 拆分了矩陣 B，將載入的資料量減半。
矩陣 B 在兩個 SM 之間共享，這意味著 B0 和 B1 塊需要透過 DSMEM 進行通訊。雖然 DSMEM 和 SMEM 之間存在頻寬差異，但由於我們載入的是較小的塊，因此對協調的影響很小。即便如此，我們懷疑在 Blackwell 上，TPC 中 SM 之間的通訊頻寬高於 DSMEM，因此 MMA.2SM 會利用這一點來實現更好的效能。
第五代 Tensor Core 除了可以執行常規矩陣乘法之外，還可以執行卷積。tcgen05.mma它支援權重平穩模式，並帶有一個收集器緩衝區，用於快取矩陣 B 以供複用。更多資訊，請參閱CUDA 文件和相應的權重平穩 MMA 指令。
在支援的資料型別方面，Blackwell 支援微縮放浮點格式 (MXFP)，包括 MXFP8、MXFP6 和 MXFP4。詳情請參閱本文。Blackwell 還支援 NVIDIA 自己的 NVFP4 格式，該格式以比 MXFP4 更精確而聞名。這可能是因為 NVFP4 的塊大小更小、縮放因子資料格式不同以及採用了兩級量化方法。
對於 Blackwell 架構，由於 FP8 和 FP6 具有相同的理論吞吐量，我們認為它們在 Tensor Core 中共享物理電路。相比之下，CDNA4 的 FP6 吞吐量是 FP8 的兩倍，因為它們的 FP6 單元與 FP4 共享資料路徑。我們認為 UDNA 架構將改為讓 FP6 單元與 FP8 共享資料路徑。
附註：結構化稀疏性
Ampere 具有 2:4 結構化稀疏性，理論上可將 Tensor Core 吞吐量提高一倍。它透過修剪權重矩陣來實現這一點，使得每 4 個元素中就有 2 個為零。在這種格式下，矩陣透過移除零元素進行壓縮，並使用額外的元資料索引矩陣記錄這些元素的位置，從而大致將記憶體使用量和頻寬減少了一半。
根據這篇來自中國工程師的微基準測試論文，Ampere 的結構化稀疏性可以在指令級實現大形狀 MMA 運算的 2 倍加速。論文還表明，在 Hopper 中，結構化稀疏性wgmma指令可以實現 2 倍加速，並節省高達 2 倍的用於載入權重的記憶體頻寬。
遺憾的是，2:4 結構化稀疏 GEMM 核心與 Hopper 上的密集 GEMM 核心相比，無法達到接近 2 倍的加速。這是因為在保持模型準確率的同時進行結構化剪枝存在困難、cuSPARSELt 核心未經最佳化以及 TDP 限制。除了中國人工智慧實驗室和少數西方實驗性研究 論文外，大多數人工智慧實驗室在生產推理中忽略了 2:4 結構化稀疏性，而專注於量化和提煉。Meta 正在 Llama 上進行實驗，但在很多情況下，這也是一條死路。
此外，目前缺乏使用 2:4 FP8 結構化稀疏性或 4:8 FP4 結構化稀疏性實現效能提升，同時保持零精度損失的封閉式或開放式模型，而且用於結構化剪枝的資源普遍匱乏。我們建議 NVIDIA 停止在主題演講和市場推廣材料中提及Jensen 數學結構化稀疏性失敗案例，除非他們開始持續展示能夠利用結構化剪枝進行推理的 SOTA 開放模型。一個好的開端是嘗試在 DeepSeek 上進行結構化稀疏性測試，並證明其效能可以疊加到其他技術之上，例如 NVFP4 的蒸餾和量化。
NVIDIA 在第五代 Tensor Core 中，為 NVFP4 資料型別引入了成對 4:8 結構化稀疏性。在此方案中，每八個元素被分組為四對連續元素，其中恰好兩對必須包含非零值，而其餘兩對則被修剪為零。由於 NVFP4 是子位元組資料型別，我們認為這一約束促使 NVIDIA 採用了成對 4:8 模式。雖然 4:8 稀疏性可能看起來比之前的 2:4 模式更寬鬆，但增加的成對要求意味著，對於尋求在修剪的同時保持模型精度的機器學習工程師來說，在實踐中，它並不是一個更寬鬆的約束。
張量核心尺寸增加
歷代以來，NVIDIA 擴充套件 Tensor Core 大小的步伐遠大於增加 Tensor Core 數量。NVIDIA 選擇擴充套件張量核心大小而非核心數量，是因為它更符合矩陣乘法的效能特徵。具體而言，當擴充套件問題規模時，矩陣乘法計算量呈立方增長，而資料移動量呈平方增長，這意味著運算強度呈線性增長。O(n) 的運算強度，加上資料移動量比計算量更昂貴的事實，刺激了張量核心大小的增加。
然而，增加核心大小和核心數量都會以犧牲量化效應為代價。具體來說，核心數量過多會受到“圖塊量化效應”的影響，而核心大小過大則會導致“波量化效應”。當工作單元數量不能被工作器數量完全整除時，就會出現“波量化效應”，導致在處理最終較小批次工作時利用率下降。增加張量核心大小本質上就是增加工作單元大小，導致小型矩陣的利用率低下。
運算強度的線性增長也推動了 MMA 形狀的增加。更大的 MMA 形狀可以增強運算元共享粒度。具體來說，啟動更少的較大塊可以提高資料重用率，從而節省記憶體佔用以及 RF 和 SMEM 的頻寬。對於 Blackwell 之前的架構，這導致執行 MMA 操作所需的執行緒數量不斷增加，從 8 個執行緒的四對 (Volta) 到 32 個執行緒的 Warp (Ampere)，再到 128 個執行緒的 Warpgroup (Hopper)。
增加記憶體大小
共享記憶體幾乎每一代都在增加，而暫存器檔案的大小卻保持不變。這是因為 Tensor Core 吞吐量的提升需要更深的暫存緩衝區。
由於 Tensor Core 的資料處理速度遠超全域性記憶體的載入速度，我們使用暫存記憶體來緩衝資料，這樣記憶體載入就可以先於 MMA 操作進行。Tensor Core 的吞吐量每一代都翻一番，但全域性記憶體的載入延遲卻非但沒有降低，反而有所增加。因此，我們需要增加暫存記憶體的大小來緩衝更多資料。為了實現這一點，NVIDIA 選擇共享記憶體作為 Tensor Core 的暫存記憶體，這也解釋了為什麼共享記憶體增加了，而暫存器檔案的大小卻保持不變。
然而，Blackwell 的共享記憶體大小與 Hopper 相比並沒有增加。這是因為 tcgen05 MMA 可以利用 2 個 SM，因此每個 SM 的共享記憶體只需載入一半的運算元。因此，Blackwell 的共享記憶體大小實際上翻了一番。
NVIDIA 的暫存記憶體選擇也解釋了運算元位置逐漸從暫存器移至共享記憶體的原因。即便如此，NVIDIA 在 Blackwell 上添加了 TMEM，以支援更高的 Tensor Core 吞吐量。由於 TMEM 的位置更靠近 Tensor Core，因此可以提高能效。此外，擁有獨立的記憶體還可以增加總記憶體頻寬，從而滿足 Tensor Core 的負載需求。
在所有運算元中，矩陣 D 始終位於 TMEM 中。由於矩陣 D 的訪問頻率高於矩陣 A 和 B，我們可以利用這種設計來提升 TMEM 的能效。例如，在一個樸素的分塊矩陣乘法中，為了計算一個分塊，矩陣 D 的分塊會被訪問 2Kt 次（Kt 次讀取和 Kt 次寫入。Kt：K 維度上的分塊數量），而矩陣 A 的分塊和矩陣 B 的分塊僅會被訪問一次。
MMA教學的非同步性
其中的“H”UTCHMMA,HGMMA,HMMA 代表半精度，因為它是16位格式；“Q”代表QGMMA,UTCQMMA四分之一精度（8位），因為8位是全精度（32位）的四分之一。“O”代表“八進位制”，即FP4中32位的八分之一UTCOMMA。
MMA 指令看似從同步跳到了非同步。實際上，由於需要重疊，MMA 指令在 SASS 級別逐漸變為非同步LDSM指示。
在 SASS 級別，MMA 操作涉及執行一條LDSM指令，將矩陣塊從共享記憶體載入到暫存器檔案，然後執行兩條HMMA執行 MMA 指令。在執行過程中，兩個HMMA指令非同步發出，並透過硬體互鎖阻止暫存器的使用。由於硬體互鎖不允許重疊的 LDSM 指令，因此順序執行一條LDSM或兩條HMMA指令在指令釋出流水線中會產生一個小氣泡。然而，隨著 Tensor Core 的速度越來越快，這個氣泡會造成不可忽略的效能損失，因此需要為 MMA 引入非同步完成機制。
Hopper 支援非同步完成機制提交和隔離wgmma。HGMMA發出指令時，沒有硬體互鎖來保護暫存器的使用。相反，編譯器會排程LDSM下一個 MMA 指令，並使用FENCE保持下一個HGMMA等待的指令。使用 Blackwell，MMA 操作是完全非同步的。載入到 Tensor Memory（tcgen05.ld / tcgen05.st / tcgen05.cp）的指令全部是顯式非同步的。
資料型別精度降低
每一代 NVIDIA Tensor Core 中，NVIDIA 都會不斷新增低精度資料型別，從 16 位到 4 位。這是因為深度學習工作負載對低精度的容忍度極高。對於推理尤其如此，推理可以使用比訓練更低的精度。低精度更節能，佔用更少的矽片面積，並實現更高的計算吞吐量。在新一代 Tensor Core 中，我們還看到 NVIDIA 移除了 FP64 支援，以便在矽片面積和功耗預算允許的情況下優先處理低精度資料型別。
有趣的是，優先順序排序也影響了整數資料型別的支援。自 Hopper 以來，INT4 資料型別已被棄用，而在 Blackwell Ultra 上，我們發現 INT8 計算吞吐量有所下降。這是由於低精度整數資料型別的普及延遲造成的。儘管 Turing 支援 INT8 和 INT4，但直到 4 年後，新的推理量化方法才能夠利用 INT4 的緊湊性來服務於 LLM。那時，NVIDIA 已經在 Hopper 上棄用了 INT4 wgmma。
參考連結
https://semianalysis.com/2025/06/23/nvidia-tensor-core-evolution-from-volta-to-blackwell/#ampere
*免責宣告：本文由作者原創。文章內容系作者個人觀點，半導體行業觀察轉載僅為了傳達一種不同的觀點，不代表半導體行業觀察對該觀點贊同或支援，如果有任何異議，歡迎聯絡半導體行業觀察。
END
今天是《半導體行業觀察》為您分享的第4074期內容，歡迎關注。
推薦閱讀
加星標⭐️第一時間看推送，小號防走丟
求點贊
求分享
求推薦
