极客工场

Appearance

共 5,251 字预计阅读约 18 分钟

Tensor Core 张量核心：一条指令完成矩阵块乘加的 GPU 专用单元 ​

更新时间：2026-09-25。前一篇 GPU 架构与 Warp 调度器 讲的是 GPU 的通用骨架：SM、warp、共享内存与延迟隐藏。本篇往 SM 内部再走一层，看 2017 年之后重新定义 GPU 算力账单的那块专用电路——Tensor Core（张量核心）。理解它需要三块前置知识：SM 与 warp 的硬件模型见 SM 与 CUDA Core 内部结构，各种低精度浮点格式的位级含义见 混合精度：FP16/BF16/FP8，大矩阵如何切块喂给片上存储见 矩阵乘与 Tiling 分块。而专用矩阵计算的另一条路线，可对照 脉动阵列。

本文要回答的问题 ​

  • Tensor Core 在 SM 里处于什么位置？它和 CUDA core 是替代还是并列关系？
  • 一条 MMA 指令具体完成多少计算？为什么「16×16×16」这个小块能换来数倍到数十倍的算力？
  • 程序员通过什么接口使用它（WMMA、mma.sync PTX）？矩阵碎片（fragment）在 32 个线程间如何分布？
  • FP16、BF16、TF32、FP8、INT8、FP4 各在哪一代引入，对应典型算力是多少？
  • 使用 Tensor Core 有哪些硬约束：矩阵形状、内存布局、为什么必须配合 shared memory 与 tiling？
  • Tensor Core 内部是脉动阵列吗？它和 TPU 式 systolic array 的异同在哪里？

一、从标量 FMA 到矩阵 MMA：一条指令的工作量 ​

先建立指令工作量的量级感。传统 CUDA core（本质是流水化的标量/向量 FMA 单元）一条指令让一个线程完成一次 FMA：d = a·b + c，记 2 个 FLOP。一个 warp 32 个线程齐步走，一条 warp 指令最多产出 32 次 FMA、64 FLOP。Tensor Core 改变的是「一条指令的粒度」：它是 SM 内一块独立的矩阵数据通路，一条 MMA（Matrix Multiply-Accumulate）指令由整个 warp 协同执行，对寄存器中的小矩阵块完成 D = A·B + C。

以第一代规格 m16n16k16、FP16 输入、FP32 累加为例：A 是 16×16、B 是 16×16，沿 K=16 维收缩，一次完成 16×16×16 = 4096 次乘加，即 8192 FLOP——而这只是一条指令（在时序上由若干拍的微操作流水完成）。同样一条 warp 指令，工作量是标量路线的 4096 次 FMA 对 32 次 FMA，约 128 倍的理论比值。实际芯片不会让频率和发射宽度按此比例缩放，但算力代差的数量级由此奠定：Volta V100 的 FP16 Tensor 算力 125 TFLOPS，而 FP32 通用算力约 15.7 TFLOPS，相差约 8 倍；到 A100，FP16 Tensor 稠密算力 312 TFLOPS 对 FP32 的 19.5 TFLOPS，相差约 16 倍。

宣传材料里常说 V100 有 640 个 Tensor Core（80 个 SM、每 SM 8 个），但这里的「个」并非独立处理器，而是共享同一条 MMA 数据通路的乘法器切片，按拍轮流为不同 warp 服务；真正该看的是每个 SM 每周期能发射的 MMA 数量及其 MAC 吞吐。还要注意 MMA 是多拍延迟指令，单个 warp 发射后不能立刻拿到结果，必须让多个 warp 的 MMA 交错发射才能把数据通路填满——延迟隐藏的原则与标量代码一致，只是等待的颗粒从单条 FMA 变成了整条矩阵指令。

关键认知是：Tensor Core 不是更快的 CUDA core，而是粒度完全不同的执行单元。CUDA core 仍在 SM 里，负责标量运算、地址计算、分支与非矩阵代码；Tensor Core 只在命中矩阵乘加指令时被唤醒。两者共享寄存器堆和 warp 调度，由编译器决定哪段计算落到哪条数据通路上。

二、D=A·B+C：矩阵分块的硬件结构 ​

下图给出一次 MMA 的操作数结构。A、B 是低精度输入碎片，C 是累加器初值，D 是输出；FP16/TF32/FP8 等低精度输入默认用 FP32 累加（INT8 用 INT32），累加精度与输入精度解耦，这是混合精度训练数值稳定的硬件基础。

chip-design/expert/tensor-core-matrix-acceleration.md D2 示意图

「碎片分布在 32 个线程的寄存器中」是第二个容易误解的点。不存在某个线程独占一块矩阵：以 m16n16k16 为例，每个 A 元素和 B 元素按硬件规定的映射切分到 warp 各线程的寄存器片（register fragment）里，累加器 D 的 256 个 FP32 元素也分散持有，平均每线程 8 个。程序员不能假设 D[i][j] 在哪个线程中，线程间的数据归属只对 MMA 指令有意义。因此使用 Tensor Core 的代码形态是：把数据按规定布局装进 fragment → 发射 MMA → 结果仍在 fragment 中，继续发射下一条 MMA 沿 K 维累加，直到整块算完再把 fragment 搬回共享内存或全局内存。

各代的基本块形状在扩大：Volta 是 m16n16k16；Ampere 的 mma.sync 支持 m16n8k16 等更灵活的形状并引入稀疏指令；Hopper 的 wgmma（warp group MMA）把协作粒度从一个 warp 扩大到 4 个 warp（128 线程），可直接对共享内存中的矩阵发起 64×N×K 的大矩阵乘，单条指令覆盖的 MAC 数量再上一个台阶。形状变大不是数字游戏：发射一条指令摊薄的译码、调度和寄存器搬运开销更多，且更大的 N 维让同一块 A 被更多 B 列复用。

三、为什么快：MAC 密度与单元内数据复用 ​

Tensor Core 的速度优势可以拆成三个来源。

第一是硅面积的用途。CUDA core 每个 FMA 都配有完整的独立调度入口：译码、操作数收集、旁路与流水线控制。Tensor Core 把控制成本从「每个 FMA 一份」压缩为「每整块矩阵一份」，省下来的面积全部换成乘法器阵列。同样的 SM 面积预算下，能同时翻转的乘法器数量高出一个数量级，峰值算力自然拉开差距。

第二是数据复用被做进了单元内部。在普通 CUDA core 上模拟 16×16×16 乘加，每个操作数要按线程各自从寄存器读出，沿 K 维的中间和还要在寄存器间显式流转。Tensor Core 的数据通路在硬件上组织成阵列式的局部连接：操作数载入一次，在单元内部沿收缩结构被多个 MAC 复用，K 维累加在本地寄存器链上完成，不回绕寄存器堆。这正是专用矩阵单元的通用思想——把「读一次、反复用」焊死在电路里。

第三是低精度的乘数效应。FP16 乘法器面积约为 FP32 的四分之一，FP8 再减半，FP4 再减半。Tensor Core 从第一代起就面向低精度设计，同样面积里塞入的 MAC 数量随位宽下降成倍增长。再叠加 2:4 结构化稀疏（Ampere 起，每连续 4 个权重至多 2 个非零，硬件跳零后算力标称翻倍），宣传册上的数字是多种机制连乘的结果。

对比维度CUDA core（标量 FMA 通路）Tensor Core（矩阵 MMA 通路）
一条 warp 指令的工作量32 次 FMA（64 FLOP）4096 次 MAC 起步（8192 FLOP，m16n16k16）
操作粒度标量/向量，每线程独立矩阵块，整个 warp（Hopper 起为 warp group）协同
典型精度FP32/FP64 全精度FP16/BF16/TF32/FP8/INT8/FP4，FP32/INT32 累加
数据复用软件显式搬运，经寄存器堆单元内阵列式局部连接，载入即复用
控制开销每 FMA 摊完整流水线控制每整块矩阵摊一次译码调度
适用计算任意标量代码、地址计算、非规则算子形状规整的 GEMM/卷积/注意力矩阵乘
A100 典型算力FP32 19.5 TFLOPSFP16 312 TF、TF32 156 TF、INT8 624 TOPS（稠密）

从能耗角度看收益更直观：先进工艺上一次片外 HBM 访问的能耗约为一次寄存器级加法的数百倍，而 MMA 把一次搬运摊到几十次乘法上，每 MAC 的平均能耗随复用程度直线下降。这也是 Tensor Core 在 roofline 图上抬高「算力平顶」的方式——它不改变 HBM 带宽斜线，只把 GEMM 的理论平顶抬到标量路线的十倍以上；至于实际运行能不能走到平顶，仍取决于 tiling 之后算子落在横轴上的有效算术强度。

需要强调，上表的倍数是峰值理论值。Tensor Code 只有在数据供给不断、形状对齐时才能兑现这些数字；喂不饱时它不比标量通路强，甚至因为碎片装箱开销而更差。

四、编程接口与数据通路：WMMA、fragment 与 shared memory tiling ​

软件栈分三层。最底层是 PTX 指令 mma.sync（Volta 时代为 mma），直接规定 m、n、k 形状、数据类型与布局；其上是 CUDA C++ 的 WMMA（Warp Matrix Multiply Accumulate）API，用 fragment 模板类封装矩阵碎片，提供 load_matrix_sync、mma_sync、store_matrix_sync 三个基本操作；最上层是 cuBLAS、CUTLASS 这类库，普通业务代码通常只调用 cuBLAS 的 GEMM 接口，由库内部选择 Tensor Core 实现。手写内核或做极致优化时才会下沉到 WMMA/CUTLASS。

chip-design/expert/tensor-core-matrix-acceleration.md D2 示意图

这条通路解释了为什么 Tensor Core 编程与 tiling 不可分割。寄存器装不下大矩阵，HBM 又喂不动阵列，于是每个 warp（或 warp group）先把 A、B 的一块沿 K 维切片搬到 shared memory，按 swizzle（交错重排）布局存放以避开 bank conflict，再逐片载入 fragment 发射 MMA；累加器在整个 K 维循环中始终留在寄存器，算完一个输出 tile 才写回。Ampere 提供 cp.async 让全局到共享的拷贝绕过寄存器并异步完成，配合双缓冲把搬运隐藏在计算之下；Hopper 进一步引入 TMA（Tensor Memory Accelerator），一条异步指令完成多维 tile 的整块搬运，并由 wgmma 直接消费共享内存中的操作数，大幅压低寄存器压力和指令数。

五、精度演进与各代算力：从 FP16 到 FP4 ​

Tensor Core 的演进史基本就是一部低精度格式扩张史。每一代都在「乘法器更小、单卡 MAC 更多」与「算法侧证明该精度够用」之间找平衡。

精度（位宽）引入代（典型卡）该精度典型 Tensor 算力主要用途
FP16（16bit）Volta，2017（V100）125 TFLOPS第一代混合精度推理/训练
INT8（8bit）+ 整数Turing，2018（T4）130 TOPS 量级推理量化
BF16（16bit）、TF32（19bit 输入/FP32 累加）Ampere，2020（A100）BF16/FP16 312 TF；TF32 156 TF；INT8 624 TOPS（稠密）BF16 成训练主力；TF32 让 FP32 代码自动加速；2:4 稀疏翻倍
FP8（E4M3/E5M2，8bit）Hopper，2022（H100）FP8 约 1979 TF；FP16 约 989 TF（稠密），稀疏再翻倍大模型训练与推理，FP32 约 67 TF
FP4（4bit）、更细 FP6/FP8 配比Blackwell，2024（B200）FP4 约 9 PFLOPS 稠密（稀疏约 18 P）；FP8 约 4.5 PF 量级万亿参数推理与低比特训练

训练侧的需求塑造了 BF16 与 FP8 的设计。FP16 只有 5 位指数，大梯度容易上溢，BF16 用 8 位指数换取与 FP32 相同的量程，尾数精度换稳定性；FP8 则分工使用两种格式，E4M3 有 4 位尾数、精度优先，承担前向，E5M2 有 5 位指数、量程优先，承担反向，并配合每 tile 的在线缩放因子把数值压进各自的表示范围。累加始终在 FP32 下完成，低比特只作用于乘法输入——精度演进从来不是简单砍位宽，而是输入格式、累加精度与缩放策略三层协同设计。

读这张表有三条纪律。其一，FP32 行是「不使用 Tensor Core（或走 TF32 之外的全精度通路）」的基准，算力代差要用同一张卡纵向比，而不是跨代乱比。其二，「稀疏」数字以硬件支持 2:4 结构稀疏为前提，稠密工作负载要对折看；Hopper/Blackwell 的 FP8/FP4 算力还依赖编译器真的生成对应指令，框架默认精度未必开启。其三，TF32 尤其容易误读：它的存储仍是 FP32，但乘法只取 10 位尾数（与 FP16 同位宽）再用 FP32 累加，A100 上默认接管大量 FP32 矩阵乘——想要严格 FP32 结果需显式关闭。各格式的位级取舍在 混合精度：FP16/BF16/FP8 中展开。

六、使用约束：形状、布局与 tiling 的硬门槛 ​

峰值算力不会自动兑现，工程上的约束有四条。

形状对齐：MMA 要求矩阵维度是基本块（16/8，Hopper 起涉及 64）的整数倍，K 维还必须按 k 块对齐。尾部不对齐的部分要么 mask 掉（浪费算力），要么回退 CUDA core 处理。矩阵越窄、batch 越小，边角浪费占比越高——batch=1 的自回归解码正是 GPU 利用率最低的场景之一。

布局受限：fragment 对行主序/列主序有明确规定，shared memory 中还要用 swizzle 布局换 bank conflict free 的访问。任意 stride 的通用矩阵不能直接喂，需要先重排，或由库在 tile 内处理。

累加器驻留：高性能 GEMM 的核心手法是让累加器整个 K 维循环不离开寄存器，这要求 tile 的 M×N 受限于每线程寄存器数，寄存器一旦 spill 到 local memory，Tensor Core 的优势会被访存延迟吃光，occupancy 与 tile 大小要联合调优。

再给两个更具体的坑。其一是 bank conflict：A100 上一块 16×16 的 FP16 tile 若朴素按行存放，warp 内多个线程会命中同一 bank，swizzle 把元素按地址异或错位重排，才能让一次 fragment 加载的访问分散到 32 个 bank。其二是寄存器预算：FP32 累加器加上双缓冲的 A/B 碎片，每线程寄存器占用常在两百个上下，tile 调大喂得饱阵列却压低 occupancy，调小则阵列空转——CUTLASS 的大量模板参数本质上就是在显式暴露这个折中，不存在无条件最优的一组取值。

软件依赖：Tensor Core 加速对程序员「透明」的部分只到库调用为止——cuBLAS 会自动选低精度与 Tensor Core 路径；自行编写内核时，误用 FP32 累加碎片、K 维不整除、忘记双缓冲都会让实际算力只有峰值的零头。健康的大模型训练 MFU（模型浮点利用率）在 30%~50% 区间，推理小 batch 则可能低得多。

七、Tensor Core 与脉动阵列：同构还是神似 ​

脉动阵列 一节讲过 TPU 式结构：PE 网格、数据沿固定方向每拍流动一步、权重驻留或输出驻留，没有逐 PE 的指令流。两者确实神似：Tensor Core 内部也是大规模乘法器阵列加局部寄存器连接，操作数载入后在阵列内被多个 MAC 复用，K 维累加走本地链路——从电路哲学看，它就是一块被包进 SIMT 外壳的矩阵阵列，NVIDIA 未公开完整微结构，但业界普遍认为其数据通路含收缩式组织。

差别在边界和编程模型。TPU 式脉动阵列以空间数据流为一等公民：编译器把整张计算图映射到 PE 网格，数据流动的「波前」就是调度本身，阵列尺寸大（常见 128×128、256×256），灵活性低但能效极致。Tensor Core 则是 warp 指令模型下的协处理单元：阵列较小且嵌套在 SM 中，每次 MMA 只处理小碎片，宏观调度仍由 warp scheduler、shared memory 和通用 CUDA core 完成；算不了的算子（softmax、layernorm、不规则索引）原地切到标量通路，无需跨芯片迁徙。一个是「阵列即芯片，数据流即程序」，一个是「通用 SIMT 核内嵌一小块阵列，指令粒度驱动」。还有一处量化差异值得点出：脉动阵列的利用率损失首先表现为空间上的 PE 空置——矩阵边长不整除阵列边长，整块 PE 就没有数据可做；Tensor Core 的利用率损失则分散在时间上——碎片不对齐时浪费的是若干条指令的发射槽，而 warp 调度器还可以插入其他 warp 的工作填补。前者的利用率由单次映射的几何形状决定，后者由 warp 混合与指令交错决定。这直接影响了两类芯片的部署手感：NPU 对 batch 和形状敏感、需要编译器深度整形，GPU 则对负载波动更宽容。

选择上，固定模式大规模推理是脉动阵列 NPU 的能效主场，训练与算子多样的工作负载则由 Tensor Core GPU 靠通用性和生态承接——现实产品也在互相借鉴：Hopper 的 wgmma/TMA 让 GPU 越来越像显式数据流机器，而 NPU 也普遍搭配向量核与通用核补齐灵活性。可以预期，「通用核内嵌多大阵列、阵列旁配多少通用逻辑」会继续是两边共同的收敛题目。

这一篇在知识体系里的位置 ​

总结 ​

Tensor Core 是 SM 内与 CUDA core 并列的专用矩阵单元：一条 MMA 指令由整个 warp 协同完成小矩阵块 D=A·B+C（m16n16k16 FP16 即 4096 次 MAC），把控制开销从每个 FMA 摊到整块矩阵、把数据复用做进单元内部阵列，再叠加 FP16/BF16/TF32/FP8/INT8/FP4 低精度与 2:4 稀疏，换来 A100 上 FP16 312 TF 对 FP32 19.5 TF 约 16 倍、H100 FP8 近 2 PF 的代差；但算力兑现要求维度对齐、fragment 布局合规、累加器驻留寄存器，并以 shared memory tiling、cp.async/TMA 双缓冲持续喂数；它与 TPU 式脉动阵列神似——内部都是乘法器阵列加局部复用——区别只在边界：脉动阵列以空间数据流为程序、追求极致能效，Tensor Core 以 warp 指令粒度嵌入通用 SIMT 核，靠灵活性与 CUDA 生态承接训练和不规则负载。