矩陣乘加
3 秒看懂
矩陣乘加（MMA，Matrix Multiply-Accumulate）是執行 D = A × B + C 的融合操作。它把一次矩陣乘法和一次加法合併為一個不可分割的原子運算，是當前所有 AI 加速器（GPU、TPU、NPU 等）的核心計算原語。深度學習 90% 以上的浮點運算都歸結為各種形狀的 MMA，從全連線層到卷積的 im2col 展開，再到 Transformer 的 QKV 投影和注意力分數計算。
3 分鐘產業解釋
在訓練和推論中，計算密度最高、功耗最大的不是獨個乘法，而是乘法後緊跟累加的模式。因此現代硬體將 MMA 固化到指令集或專用資料通路里，一條指令完成一個分塊矩陣的乘加，大幅減少取指、解碼、暫存器搬移開銷。
以 NVIDIA GPU 的 Tensor Core 為例，它執行固定形狀的小矩陣乘加指令（如 FP16 下 mma.sync 形狀多為 M=16, N=8, K=16），通過多次 K 維度迴圈完成分塊矩陣乘加，輸入多為 FP16/BF16/INT8，內部以更高精度（如 FP32）累加，再按需寫回較低精度結果。這種模式使具體晶片的紙面算力被標為“Tensor TFLOPS”或“TOPS”——本質上就是 MMA 指令的峰值吞吐量。
產業玩家（NVIDIA、AMD、Intel、Google、Apple、中國廠商）圍繞 MMA 在製程、頻寬、量化位寬和軟體棧上激烈競爭，因為這直接決定模型訓練速度與雲端端推論成本。
15 分鐘專家深入
MMA 在深度學習中的對映
幾乎所有神經網路核心計算均可表示為 MMA：
- 全連線層：輸入
X \in mathbb(R)^{B \times M}，權重W \in mathbb(R)^{M \times N}，直接Y = X × W。 - 卷積層：通過 im2col 或 Winograd 變換將滑動視窗乘法轉換為批次的矩陣乘。
- 多頭注意力：Q、K、V 的線性投影是獨立的 MMA；注意力權重
softmax(Q×K^T)中間也需要矩陣乘；最終加權求和A × V也是 MMA。 - RNN/LSTM 門計算：將多路輸入特徵拼接後進行矩陣乘。
- MoE（混合專家）：gate 投影、專家 FFN 全為 MMA，且 All-to-All 通訊常與 MMA 流水重疊。
硬體實現模式
- 脈動陣列（Systolic Array）：權重和輸入資料在二維陣列中按節拍流動，每個處理單元（PE）做乘加並傳遞資料。TPU v1 即以此架構實現 8-bit MMA，獲得極高能效。
- Tensor Core 類：固定尺寸的分塊矩陣乘加單元，排程器負責將大矩陣切割，填充到指令所要求的形狀。這類單元通常與 CUDA/SIMT 核心共享快取和暫存器堆。
- SIMD + 融合指令：Intel AMX 擴充套件引入獨立於 AVX-512 的二維暫存器（tile），執行
tdpbf16ps tile0, tile1, tile2之類指令，實質是對 16×16 或類似分塊做 MMA。 - NPU SoC 上的專用引擎：Apple ANE、高通 Hexagon、華為達芬奇等均包含多維 MAC 陣列，直接執行卷積或全連線層的提前量化版 MMA。
混合精度的關鍵
當前標準方案：
- 輸入與權重：FP16、BF16、INT8 或 FP8（E4M3、E5M2）。
- 累加器：FP32 或 INT32，避免中途溢位/舍入誤差。
- 輸出寫回：可保持低精度以節約頻寬。
一條 MMA 指令內部可等效為數十次常規乘加，但僅需一次解碼、一次暫存器檔案訪問。深度學習架構通過底層庫（cuBLAS、oneDNN 等）自動將高層運算元對映為最優 MMA 指令序列。
技術原理（最深，帶機制與關鍵引數）
分塊MMA: C_block += A_block × B_block
┌─────────────────────────────────────┐
│ A [M_sub × K_sub] │
│ × │
│ B [K_sub × N_sub] │
│ + │
│ C [M_sub × N_sub] ──> C_new │
└─────────────────────────────────────┘1. 指令層面
MMA 指令格式通常指定：
- 輸入矩陣 A、B 的基地址與形狀（硬體固定或引數化）
- 累加器 C 的基地址（同時為輸出 D）
- 精度格式（元素型別和累加器型別）
執行時：
- 從暫存器/共享記憶體載入 A、B 分塊到 MAC 陣列。
- 陣列內每個 PE 獨立做乘加：
c_ij += a_ik * b_kj（k 維度由陣列內部時序展開）。 - 完成 K 維度迴圈後，將 C 分塊寫回。
由於載入和計算流水重疊，隱藏延遲至關重要。輝達在 Hopper 架構引入 非同步複製 和 Warp Group 級 MMA，試圖將資料搬運完全隱藏於計算之後。
2. 計算密度與理論峰值
理論峰值 FLOPS ≈ MMA 單元數量 × 每單元每週期運算元 × 頻率。 以每週期每個 Tensor Core 完成的乘積-累加次數計為 FMA 數（1 次 = 1 乘 + 1 加，部分廠商算作 2 次操作）。深度學習中習慣用 TFLOPS/TOPS 表達。
真峰值利用率 受制於：
- 矩陣分塊是否能填滿硬體固定尺寸（tile 量化效應）
- 記憶體/快取頻寬能否餵飽 MAC 陣列（Roofline 模型中的計算/訪存比）
3. 關鍵量化引數
- 分塊尺寸：16×16、32×32、64×64 等，越大吞吐越高但暫存器壓力更大。
- K 維度步進量：一次 MMA 指令能處理的 K 維度長度（如 Pipelining K=4,8）。K 越大，計算密度越逼近理論值，但中間累加器精度要求更高。
- 資料重用：A 分塊廣播到同行所有 PE，B 分塊廣播到同列，每載入一個元素可被多次使用，這是 MMA 能效遠超標量的根本。
4. 與稀疏性結合
硬體 MMA 支援的結構化稀疏（如 2:4 模式）會跳過一半權重，要求矩陣 B 滿足特定零分佈。此時 MMA 吞吐理論上翻倍，但受開銷影響實際加速比在 1.5× 左右（[NVIDIA 官方指南]）。
技術演進史
| 時期 | 關鍵進展 | 典型硬體 / 精度 |
|---|---|---|
| 2017 前 | 標量 FMA 為主，純 SIMD 極少融合成矩陣指令。卷積通常展開為低效的標量運算或 cuBLAS GEMM | Pascal GPU（無專用 MMA 指令） |
| 2017 | Volta 推出第一代 Tensor Core：FP16 MMA，引入 FP16 乘、FP32 累加模式。每個 Tensor Core 每時鐘執行 4×4×4 的乘加操作，或一條 WMMA 指令可完成諸如 16×16×16 的分塊乘加。 | Tesla V100, 峰值 ~125 Tensor TFLOPS |
| 2018-2019 | Turing 擴充套件 INT8、INT4 等混合精度；支援更大尺寸的矩陣乘加（如 16×8×16 的 FP16 MMA） | T4 |
| 2020-2021 | Ampere 架構引入 TF32、BF16、結構化稀疏，BF16 成為訓練標準；同時脈動陣列商品化：Google TPU v3/v4，Apple M1 ANE，Intel AMX。硬體原生支援矩陣指令 | A100, Sapphire Rapids（AMX-BF16），M1/M2 |
| 2022-2023 | Hopper (H100) 引入 FP8（E4M3、E5M2）和 Warp Group MMA。非同步執行、TMA 單元專為 MMA 搬運資料。 | H100 (FP8 稠密峰值 ~1979 TFLOPS，稀疏標稱 ~3958 TFLOPS) |
| 2024+ | 加速器軟體棧深度適配 Transformer：FlashAttention-2/3 利用 MMA 指令手動排程 tiling 和重計算，最大化計算密度。 | H200, B200, Intel Gaudi 3 等 |
注：具體 FLOPS 與分塊尺寸來自廠商公開資料，實際吞吐與散熱及負載相關。
技術路線對比
| 維度 | GPU Tensor Core | 專用脈動陣列 (TPU) | CPU 矩陣引擎 (AMX) | 移動/邊緣 NPU |
|---|---|---|---|---|
| 指令集介面 | 緊密整合在 CUDA 執行緒內，Warp 級協作 | 固化成 DSA，通過架構編譯器排程 | x86 指令級 tdpbf16ps 等，tile 架構 | 多為黑盒 API，交編譯器/Middleware 對映 |
| 精度支援 | 寬泛，FP64 至 INT4，逐步過渡 FP8 | 以 BF16/INT8 為主，TPU v5p 加入 FP8 | AMX: BF16, INT8；後續可能擴充套件 FP8 | INT8 為主，部分支援 FP16 |
| 矩陣形狀靈活性 | 中等（固定分塊），通過驅動庫適配任意矩陣尺寸 | 較低，通常要求某個維度對齊極大值（如 128 或 256）才能高效 | 中等，tile 暫存器可動態配置 | 低，通常以圖層（Conv/FC）為粒度 |
| 生態成熟度 | 極高，cuBLAS / CUTLASS / PyTorch 原生 | 高，JAX/TF/PyTorch-XLA 深度最佳化 | 中，oneDNN/OpenVINO 生態快速追趕 | 碎片化，各自 SDK |
| 能效 （TOPS/W）定性 | 中等（需維持龐大快取和排程開銷） | 高（脈動結構天然低功耗） | 中高 | 極高（專為低功耗設計） |
上下游
- 上游：代工廠先進製程（決定電晶體密度、SRAM 容量）、HBM/先進封裝（CoWoS 等）、EDA/IP 供應商（提供 MAC 陣列設計庫）
- 中游：AI 晶片設計公司（自研 MMA 單元）、AI 伺服器/板卡 ODM
- 下游：深度學習架構（PyTorch, JAX, TensorFlow）、底層計算庫（cuBLAS, CUTLASS, oneDNN）、雲端服務商、模型開發商
MMA 作為中心運算元，向下驅動視訊記憶體頻寬、NoC 拓撲設計；向上定義軟體排程粒度和編譯最佳化策略。
關鍵指標
| 指標 | 說明 | 典型值範圍 [估算或定性] |
|---|---|---|
| 理論 Tensor FLOPS | 單位時間可執行的 MMA 乘加操作次數 | 高階 GPU 數百至數千 TFLOPS；移動 NPU 數至數十 TOPS |
| MMA 利用率 | 實際吞吐 / 理論峰值 | 大矩陣 GEMM 能達到 80‑95%；注意力等不規則形狀可能跌至 30‑50% |
| 算術強度 | 每次訪存可完成的浮點運算數 | Transformer 訓練中 MMA 部分的算術強度極高（數百 FLOP/Byte），而逐元素操作很低 |
| 能效比 | FLOPS / W | 雲端端推論晶片常以 1‑2 TFLOPS/W 為目標（BF16） |
| 精度轉換開銷 | 混合精度帶來的精度損失 | 採用 FP32 Accum 的 BF16 MMA 幾乎無損；INT8 量化需校準，精度損失可控制在 1% 以內（[相關文獻]） |
具體數值依賴微架構與時鐘，此處僅為量級參考。
供需與市場資料
（以下資料來自行業估算和部分廠商財報，無進一步權威來源時標為 估算）
- GPU/加速器級：2024 年資料中心 AI 晶片出貨量中，支援高吞吐 MMA 的晶片（H100/H200/ B200 等）佔比超過 85% [Jon Peddie Research 估算]。MMA 硬體供給被先進封裝（CoWoS）產能嚴重製約。
- CPU 級：AMX 的滲透率隨 Sapphire Rapids/EMR 伺服器更新而快速提升，預計 2025 年 60% 以上的新部署雲端例項將具備 BF16 AMX 能力。
- 端側：AI PC 和 AI 手機對 NPU MMA 單元的需求驅動高通驍龍 8 Gen 4、Apple A17 Pro / M4 等大幅提高 MAC 陣列規格。據 Counterpoint 估算，2024 年具備專用 MMA 硬體的手機晶片佔比已達 50%。
- 成本側：MMA 單元面積通常佔 AI 加速器 die 面積的 30‑40%，HBM 介面和快取佔其餘大部。先進製程一個 8-bit MAC 單元的面積正隨製程微縮持續下降，但 SRAM 縮放放緩，導致片上快取更趨珍貴。
代表公司與資本對映
| 環節 | 公司 | MMA 相關核心業務 / 產品 |
|---|---|---|
| GPU/加速器 | NVIDIA | Tensor Core（從 Volta 到 Blackwell 持續迭代） |
| 專用 ASIC | TPU（脈動陣列 MMA） | |
| CPU 廠商 | Intel, AMD | Intel AMX, AMD 未來 matrix 擴充套件 |
| 移動 NPU | Apple, 高通, 聯發科, 三星 | ANE / Hexagon / 聯發科 APU / 三星 NPU |
| 中國 AI 晶片 | 華為, 寒武紀, 海光, 燧原等 | 各自研發的達芬奇/智慧處理器核心包含 MMA 單元 |
| 軟體生態 | 輝達 (CUTLASS), Intel (oneDNN) | 提供 MMA 模板庫，使上層架構高效對映 |
| 代工與封裝 | 台積電, 三星, Intel | 高階 MMA 晶片不可或缺的製程與 CoWoS 產能 |
資本對映：
- 硬體提供 → 先進製程 + 封裝 + HBM 供應量決定 MMA 產能，相應 ETF/個股為市場焦點。
- 軟體競爭力 → 若公司僅擁有 MMA 硬體但編譯器/庫滯後，實際利用率將遠低於競品，形成投資陷阱。
- FP8/稀疏化 → 新一代 MMA 對 FP8 支援程度直接影響下一代大型模型訓練成本，成為估值分水嶺。
投資邏輯
- 供需瓶頸：MMA 晶片的核心瓶頸不在 MMA 單元本身，而在 HBM 和先進封裝。追蹤 CoWoS 產能和 HBM 交期更能預判 AI 算力供給。
- 混合精度代際紅利：FP8 的 MMA 比 BF16 能效提升約 2×（同等矩陣規模下），率先穩定部署 FP8 的硬體將搶下訓練市場的大份額。
- 軟體棧護城河：MMA 利用率高度依賴底層庫（如 CUTLASS 最佳化的塊大小、流水編排）。代差不僅體現在硬體 FLOPS，更體現在軟體利用率“天花板”。
- 端側滲透：當邊緣 NPU 的 MMA 效能超過 10 TOPS（INT8） 且功耗 <1W 時，裝置端執行 7B 引數模型成為現實，重構移動 SoC 估值體系。
- 風險因子：稀疏 MMA 收益未達預期、結構定製過度（無法適配新型模型運算元）、能效改善不及製程提升成本，均可能導致某些架構被邊緣化。
常見誤讀糾偏
誤讀 1：“一臺 GPU 的 TFLOPS 就是實際訓練速度”
事實：GPU 紙面 TFLOPS 是 MMA 單元 100% 佔用的理論值。實際訓練中，受限於視訊記憶體頻寬、核心啟動開銷、非 MMA 模組（Softmax、LayerNorm）及通訊延遲，端到端利用率通常在 40‑60%（大規模並行訓練中可能更低）。投資者若僅根據紙面算力橫向對比硬體，會嚴重誤判價效比。
誤讀 2：“MoE 模型的引數總量除以 expert 數就是每個 token 啟用的 MMA 計算量”
事實：MoE 模型中單個 token 的路由僅啟用 1‑3 個專家，但共享層（自注意力等）仍為全部引數載入並計算。啟用引數量 ≠ 總引數量 / expert 數，而是共享層引數 + 少量專家引數。對應的 MMA 計算量遠小於總引數的 MMA 量，直接用總引數量除 expert 數估算 FLOPS 是常見技術錯誤。
學習路徑
- 基礎：理解矩陣乘法與累加的數學定義，學習 GEMM 的分塊演算法（tiling）。
- 硬體入門：研讀 NVIDIA Tensor Core 白皮書或 Google TPU 架構簡述，理解脈動陣列與暫存器級 MMA 的區別。
- 效能模型：用 Roofline 模型分析計算密集 vs 訪存密集的界限，嘗試用
ncu/nvvp分析 MMA 利用率。 - 軟體庫：閱讀 CUTLASS 原始碼的 GEMM 示例，關注
wmma和mmaPTX 指令的使用與流水編排。 - 前沿最佳化：精讀 FlashAttention-2/3 論文的 MMA 排程、非同步複製和重計算策略，理解如何將不規則注意力對映為 MMA 友好形式。
一句話總結
矩陣乘加是 AI 算力的“原子操作”，任何脫離 MMA 去討論深度學習晶片效能與投資價值的分析，都不觸及數字背後的物理本質。
延伸閱讀與來源
- NVIDIA A100/H100 白皮書：Tensor Core 架構說明
- CUTLASS 開源庫：GitHub - NVIDIA/cutlass
- Google TPU 論文：Jouppi et al., “In-Datacenter Performance Analysis of a Tensor Processing Unit”, ISCA 2017
- Intel AMX 手冊：Intel Architecture Instruction Set Extensions
- FlashAttention-3：Dao et al., “FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision”, arXiv 2024
（以上來源均為行業公認的公開技術文件或學術論文，具體規格標註以廠商釋出物為準。搜尋失敗，未獲得第三方獨立供應鏈資料。）