矩阵乘加

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：

1. 全连接层：输入 X \in mathbb(R)^{B \times M}，权重 W \in mathbb(R)^{M \times N}，直接 Y = X × W。
2. 卷积层：通过 im2col 或 Winograd 变换将滑动窗口乘法转换为批量的矩阵乘。
3. 多头注意力：Q、K、V 的线性投影是独立的 MMA；注意力权重 softmax(Q×K^T) 中间也需要矩阵乘；最终加权求和 A × V 也是 MMA。
4. RNN/LSTM 门计算：将多路输入特征拼接后进行矩阵乘。
5. 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）
• 精度格式（元素类型和累加器类型）

执行时：

1. 从寄存器/共享内存加载 A、B 分块到 MAC 阵列。
2. 阵列内每个 PE 独立做乘加：c_ij += a_ik * b_kj（k 维度由阵列内部时序展开）。
3. 完成 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 官方指南]）。

技术演进史

注：具体 FLOPS 与分块尺寸来自厂商公开数据，实际吞吐与散热及负载相关。

技术路线对比

上下游

• 上游：代工厂先进制程（决定晶体管密度、SRAM 容量）、HBM/先进封装（CoWoS 等）、EDA/IP 供应商（提供 MAC 阵列设计库）
• 中游：AI 芯片设计公司（自研 MMA 单元）、AI 服务器/板卡 ODM
• 下游：深度学习框架（PyTorch, JAX, TensorFlow）、底层计算库（cuBLAS, CUTLASS, oneDNN）、云服务商、模型开发商

MMA 作为中心算子，向下驱动显存带宽、NoC 拓扑设计；向上定义软件调度粒度和编译优化策略。

关键指标

具体数值依赖微架构与时钟，此处仅为量级参考。

供需与市场数据

（以下数据来自行业估算和部分厂商财报，无进一步权威来源时标为 估算）

• 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、苹果 A17 Pro / M4 等大幅提高 MAC 阵列规格。据 Counterpoint 估算，2024 年具备专用 MMA 硬件的手机芯片占比已达 50%。
• 成本侧：MMA 单元面积通常占 AI 加速器 die 面积的 30‑40%，HBM 接口和缓存占其余大部。先进制程一个 8-bit MAC 单元的面积正随制程微缩持续下降，但 SRAM 缩放放缓，导致片上缓存更趋珍贵。

代表公司与资本映射

资本映射：

• 硬件提供 → 先进制程 + 封装 + HBM 供应量决定 MMA 产能，相应 ETF/个股为市场焦点。
• 软件竞争力 → 若公司仅拥有 MMA 硬件但编译器/库滞后，实际利用率将远低于竞品，形成投资陷阱。
• FP8/稀疏化 → 新一代 MMA 对 FP8 支持程度直接影响下一代大模型训练成本，成为估值分水岭。

投资逻辑

1. 供需瓶颈：MMA 芯片的核心瓶颈不在 MMA 单元本身，而在 HBM 和先进封装。跟踪 CoWoS 产能和 HBM 交期更能预判 AI 算力供给。
2. 混合精度代际红利：FP8 的 MMA 比 BF16 能效提升约 2×（同等矩阵规模下），率先稳定部署 FP8 的硬件将抢下训练市场的大份额。
3. 软件栈护城河：MMA 利用率高度依赖底层库（如 CUTLASS 优化的块大小、流水编排）。代差不仅体现在硬件 FLOPS，更体现在软件利用率“天花板”。
4. 端侧渗透：当边缘 NPU 的 MMA 性能超过 10 TOPS（INT8） 且功耗 <1W 时，设备端运行 7B 参数模型成为现实，重构移动 SoC 估值体系。
5. 风险因子：稀疏 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 是常见技术错误。

学习路径

1. 基础：理解矩阵乘法与累加的数学定义，学习 GEMM 的分块算法（tiling）。
2. 硬件入门：研读 NVIDIA Tensor Core 白皮书或 Google TPU 架构简述，理解脉动阵列与寄存器级 MMA 的区别。
3. 性能模型：用 Roofline 模型分析计算密集 vs 访存密集的界限，尝试用 ncu/nvvp 分析 MMA 利用率。
4. 软件库：阅读 CUTLASS 源码的 GEMM 示例，关注 wmma 和 mma PTX 指令的使用与流水编排。
5. 前沿优化：精读 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

（以上来源均为行业公认的公开技术文档或学术论文，具体规格标注以厂商发布物为准。搜索失败，未获得第三方独立供应链数据。）