香港飛龍.online 官方授權發布的第4代「香港飛龍」標誌 本文内容： 公衆號記得加星標??，第一時間看推送不會錯過。來源：內容編譯自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 的文檔並不完善。VoltaNVIDIA 爲何添加 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 MMATesla 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 的混合精度訓練技術相呼應。該技術表明，可以在不損失模型精度的情況下以較低的精度訓練模型。TuringTuring 架構包含第二代 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級異步 MMANVIDIA 通過 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 第五代 MMATensor 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*免責聲明：本文由作者原創。文章內容系作者個人觀點，半導體行業觀察轉載僅爲了傳達一種不同的觀點，不代表半導體行業觀察對該觀點贊同或支持，如果有任何異議，歡迎聯繫半導體行業觀察。今天是《半導體行業觀察》爲您分享的第4074期內容，歡迎關注。加星標??第一時間看推送，小號防走丟求推薦 (本文内容不代表本站观点。) --------------------------------- |
本網站是"非商業"(non-commercial)。 暫統計至2025年6月25日，本網站共113091篇文章。 本網頁使用CDN cache緩存，顯示的網頁內容可能並非最新版本。
