张量指令 ISA：从标量 FMA 到一条 mma 指令上千次矩阵乘加
更新时间：2026-09-25。姊妹篇 Tensor Core 张量核心 从硬件单元视角讲了 SM 内那块矩阵电路，本篇换一个视角：站在指令集架构（ISA, Instruction Set Architecture）一侧，看指令本身的语义如何从「一个周期一算」一步步走到「一条指令一个小矩阵」。理解本文需要的基础：乘加运算的电路原子见 MAC 乘加单元，SIMD 的编程模型可对照 OpenCL 与 SYCL 总览，GPU 的 warp 执行模型见 GPU 架构与 Warp 调度器。
本文要回答的问题
- 标量指令、SIMD/向量指令、张量指令三者的粒度差多少？一条指令各完成多少 MAC？
- 为什么 ISA 的语义要向矩阵倾斜？译码与控制开销是怎么被摊薄的？
- 张量指令的操作数放在哪里？寄存器为什么要按 tile 而不是按标量组织？
- WMMA、
mma.sync、wgmma三层接口谁高谁低，程序员实际接触哪一层？ - 编译器和汇编器如何把高级代码变成一条矩阵指令？写 PTX 内联汇编是什么形态？
- ISA 与微架构如何互相锁死？x86 AMX、ARM SME 又给出了什么不同答案？
一、指令粒度的三级演进
指令集发展史可以用「一条指令携带多少计算」来重新分期。
第一级是标量指令。传统 RISC/GPU 通用通路的 FMA 语义是 d = a*b+c，操作数是三个寄存器，一条指令 1 次乘加、2 FLOP。处理器内部的取指、译码、发射、旁路控制都为这一次计算服务，控制逻辑与运算逻辑的比例接近一比一。
第二级是 SIMD（Single Instruction Multiple Data，单指令多数据）与向量指令。x86 的 AVX-512 把寄存器扩成 512 位，一条 vfmadd 同时对 16 个 FP32 lane 做 FMA，即 16 次乘加、32 FLOP；ARM NEON 是 4 个 FP32 lane，RISC-V V extension 的 lane 数由实现决定。注意 SIMD 改变的是宽度而非语义：指令描述的仍是「对一批互不相关的标量分别做同一个运算」，lane 与 lane 之间没有数据流动。
第三级就是张量指令（tensor instruction）。NVIDIA PTX 的 mma.sync.m16n8k16.f16 语义是：取 16×16 的 A 碎片与 16×8 的 B 碎片，沿 K=16 维收缩，累加到 16×8 的 C 上，一次完成 16×8×16 = 2048 次 MAC（4096 FLOP）。指令语义第一次不再是「一批独立标量运算」，而是一个有内禀归约结构的矩阵算子：输出元素之间共享输入，K 维必须在指令内部完成求和。到 Hopper 的 wgmma，协作范围扩到 128 线程的 warp group，FP8 下单条指令可达 m64n256k16，即 64×256×16 = 262144 次 MAC——一条指令承载的计算量是标量 FMA 的 26 万倍。
三级之间不是替代而是叠加：今天的 GPU 同时保留标量通路（地址计算、分支）、SIMD 式向量化（elementwise 算子）和张量通路（矩阵乘），编译器按代码形态选择落点。
二、指令语义为什么向矩阵倾斜
把矩阵乘加做成一条独立指令，动力来自一笔控制开销账。
处理器每执行一条指令，都要支付一笔与运算本身无关的固定成本：取指带宽、译码、重命名、发射仲裁、操作数收集、完成提交。对标量 FMA，这笔成本由 1 次 MAC 承担；对 AVX-512，由 16 次 MAC 分摊；对 mma.sync.m16n8k16，由 2048 次 MAC 分摊。每 MAC 的平均控制成本降到两百分之一以下，省下来的功耗与硅面积可以换成更多乘法器——这正是 Tensor Core 张量核心 中峰值算力跃迁的 ISA 侧解释。反过来看，矩阵计算也恰好适合被「一条指令化」：它形状规整、无分支、数据依赖只有沿 K 维的归约，是固定成本摊销最理想的载体。
| 对比维度 | 标量指令（FMA） | SIMD/向量指令（vfmadd） | 张量指令（mma/wgmma） |
|---|---|---|---|
| 操作数形态 | 3 个标量寄存器 | 2~3 个宽向量寄存器 | 2~3 个矩阵 tile（碎片跨线程分布） |
| 每指令 MAC 数 | 1 | 16（AVX-512 FP32） | 2048（m16n8k16）；wgmma 至多 262144 |
| 每指令 FLOP | 2 | 32 | 4096；wgmma FP8 至 524288 |
| lane/元素间关系 | 无 | 互不相关 | K 维归约、输出共享输入 |
| 译码/控制开销摊到 | 1 个 MAC | 16 个 MAC | 2048 个以上 MAC |
| 典型用途 | 地址计算、分支、标量代码 | elementwise、归一化 | GEMM、卷积、注意力矩阵乘 |
第二层收益是指令带宽。若用标量指令完成 m16n8k16 这一块，需要发射 2048 条 FMA 外加全部地址与循环控制指令；张量指令把它压成一条，前端取指与发射带宽几乎不再构成瓶颈。这也解释了为什么粒度演进没有停在 SIMD：SIMD 只摊薄了数据通路宽度，指令仍按拍逐条发射，而矩阵乘的 K 维循环本身是可压缩的固定模式——干脆把循环也收进指令语义里。
那为什么不把指令做得无限大，一条指令算完整张矩阵？粒度选择受三个反作用力约束。其一是延迟：一条 mma 指令本身是多拍操作，粒度越大、单条延迟越长，单个 warp 必须等待，需要更多 warp 交错发射才能填满通路；其二是寄存器压力：tile 越大，驻留的碎片元素越多，能同时驻留的 warp 数（occupancy）反而下降；其三是尾部浪费：形状对齐是硬契约，指令越大，非整除矩阵和小 batch 时整块空置的代价越高。从 16×16 到 64×256 的每一步放大，都是在这三个约束重新放松（更多 warp、更大共享内存、TMA 自动搬运）之后才成立的。
三、寄存器按 tile 组织：操作数不再是标量
指令语义放大后，第一个被改造的就是寄存器模型。标量 ISA 的寄存器是一维标量槽；SIMD 把槽加宽；张量指令则要求寄存器（或寄存器碎片的集合）能表示二维 tile。
NVIDIA 路线没有引入独立的矩阵寄存器，而是沿用 SIMT 模型：tile 被切碎分布在 warp 32 个线程的普通寄存器中。以 m16n8k16 FP16 为例，A 碎片 16×16、B 碎片 16×8，按 PTX 规定的映射切到各线程，累加器 C 的 128 个 FP32 元素平均每线程持有 4 个。线程私有的只是若干元素，矩阵只在「整个 warp 共同发射一条 mma」时作为整体存在。这种做法的好处是零新增架构状态、碎片还能被普通 load 指令填充；代价是映射规则复杂，程序员必须按布局装箱。
碎片映射还带着布局语义。助记符里的 .row.col 后缀分别声明 A、B 是行主序还是列主序，选错后缀不会报错却会算出错误结果；从共享内存装填碎片时还要配合 swizzle 交错布局避开 bank conflict。也就是说，装箱责任在软件，正确性也在软件。
其他架构选择了显式 tile 寄存器，形成值得对照的两条路线。Intel AMX（Advanced Matrix Extensions）引入 8 个 TMM tile 寄存器，每个可达 1KB（16 行 × 64 字节），使用前要先用 ldtilecfg 配置每个 tile 的行列数（palette 配置本身是一块架构状态），之后 tdpbf16ps %tmm1, %tmm2, %tmm0 这样的指令直接以 tile 名为操作数；ARM SME（Scalable Matrix Extension）则提供一块可按二维切片寻址的 ZA 存储，矩阵外积指令把 tile 作为一等公民，并随流式向量长度（SVL）伸缩。它们与 NVIDIA 方案回答的是同一个问题——矩阵操作数必须有一块按二维组织、供阵列直接消费的架构状态——区别只在状态挂在线程私有 tile 上，还是挂在 warp 协作的寄存器碎片集合上。AMX/SME 路径的映射直观、单线程即可发起，代价是新增的架构状态需要显式配置与保存恢复；碎片路径零新状态、复用现有 load 通路，代价是映射规则隐晦、必须凑齐整个 warp。
tile 化寄存器还带来一条硬约束：形状被编码进架构。m、n、k 不是运行时参数而是指令助记符的一部分，硬件只认少数几种尺寸，矩阵维度必须是基本块的整数倍。ISA 在获得摊销收益的同时，把「只服务规整矩阵」写成了契约。
四、暴露给软件：WMMA → mma.sync → wgmma，越来越底层
同一个硬件能力，NVIDIA 在软件栈上开了三个高度不同的窗口，演进方向是一层层往下放权。
WMMA 是最厚的一层：CUDA C++ 用 fragment<matrix_a, 16, 16, 16, half> 描述碎片，load_matrix_sync / mma_sync / store_matrix_sync 三个调用包办装箱、发射与取回，线程到元素的映射完全藏在模板里。再往下是 PTX 的 mma.sync：形状、精度、行主序/列主序全部显式写出，sync 表示指令在 warp 汇合点语义下同步完成，发射后结果即在寄存器碎片中可见。
Hopper 的 wgmma 把抽象层再削薄一层。它允许 128 线程的 warp group 直接对 shared memory 中的矩阵发起 64×N×K 乘加，操作数不必先装进每线程寄存器；更关键的是它变成异步指令，软件必须显式管理三条配套指令：wgmma.fence 声明共享内存可见性、wgmma.commit_group 提交一组发射、wgmma.wait_group 等待完成。从 WMMA 到 wgmma 的方向很清楚：ISA 把越来越多的调度与数据排布决策权交还软件，以指令数和同步复杂度换取对阵列喂数节奏的完全控制。CUTLASS 3.x 正是围绕 wgmma/TMA 重写的。
五、编译器与汇编如何生成张量指令
软件到指令的下降（lowering）有三条典型路径。
路径一是库路径，绝大多数应用走这里：框架调用 cuBLAS 的 GEMM，库内部以 CUTLASS 风格的 C++ 模板实例化出特定形状、精度的内核，nvcc 编译模板时把 mma_sync 内部函数下降为 PTX mma，再由 ptxas 汇编成 SASS 机器码（SASS 中对应助记符为 HMMA/MUFU 一类）。形状在编译期固化，这与第三节「形状写进 ISA」互为表里。
路径二是手写下降。极致优化的内核会直接在 CUDA 里内联 PTX，形如 asm volatile("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 ...")，操作数顺序、布局后缀、寄存器约束全部手工给出；wgmma 内核还要手写 fence/commit/wait 的节拍。这类代码通常用内联函数封装成小原语，再由上层模板批量组合。
路径三是深度学习编译器：Triton、XLA、TVM 这类系统从块级 IR 出发，自动完成 tile 划分、共享内存排布与双缓冲，后端直接生成 mma.sync 或 wgmma 指令；跨架构侧，LLVM/GCC 为 AMX 提供内置函数、为 SME 提供 arm_sme 内联函数层。无论哪条路径，共同的结构都是「高层描述块矩阵运算 → 编译期选定形状与布局 → 下降到固定编码的单条矩阵指令」。程序员能在任意一层介入，但越往下越要亲自承担碎片映射与同步正确性。
一个细节值得点出：K 维循环在三层路径里都不被收进单条指令——每条 mma 只推进一个 k 块，累加器碎片留在寄存器中，由软件循环反复发射 mma 沿 K 维叠加，直到整块完成。编译器真正做的决策是：输出 tile 取多大、K 每次推进多少、几条 mma 交错、搬运指令插在哪两个发射之间。这解释了为什么同一段 GEMM 在不同内核里性能可以相差一倍以上——ISA 只提供固定形状的积木，拼积木的手艺仍在软件。
六、架构与 ISA 的耦合：契约如何锁死双方
张量指令把 ISA 与微架构的耦合提到了少见的紧度。
对硬件，ISA 契约规定了阵列的最小服务形状、操作数映射与累加精度，数据通路必须按此组织；但契约也留出自由：同一条 mma.sync.m16n8k16 在 Volta、Ampere 上可以有完全不同的延迟和吞吐实现，软件看到的语义不变。对软件，耦合则表现为一组可见的代际特征：结构化稀疏在 Ampere 上以专门的稀疏 mma 指令暴露，2:4 压缩模式从此写进指令编码；Hopper 必须新增 TMA 与 wgmma 三条同步指令才能放出共享内存直连的新数据通路；Blackwell 的 FP4/FP6 同样以新助记符进入 PTX。每一次微架构的算力跃迁，都伴随 ISA 扩口——没有指令暴露，再多乘法器软件也用不到。
这种紧耦合也给出了生态分野。NVIDIA 用「固定小形状 + warp 碎片映射 + 逐层放权」保持 PTX 的稳定与编译便利；AMX 用显式 1KB tile 寄存器换简洁语义，但配置 tile 尺寸需要额外的 ldtilecfg 状态管理；SME 沿用可伸缩向量长度哲学，tile 切片随实现缩放。三种答案没有绝对优劣，取舍都在同一条轴上：把多少矩阵结构显式写进架构状态与指令编码，把多少调度自由留给软件。判断一个张量 ISA 设计的好坏，最终就看两件事——固定控制成本被摊到了多少 MAC 上，以及软件为吃到这个摊销要付出多少装箱、对齐与同步代价。
这一篇在知识体系里的位置
- 硬件单元侧：Tensor Core 张量核心 讲同一对象的微结构与算力数字；
- 电路原子：MAC 乘加单元 解释被张量指令批量驱动的基本部件；
- 并行编程参照：OpenCL 与 SYCL 总览 对照 SIMD/向量模型；
- 执行环境：GPU 架构与 Warp 调度器 是 warp 协作发射的大背景；
- 另一条专用路线：AI 加速器脉动阵列 对比「数据流即程序」与「指令粒度驱动」。
总结
张量指令 ISA 的本质是把指令粒度从标量（1 MAC/指令）、SIMD（16 MAC/指令，lane 互不相关）推到矩阵 tile（mma.sync.m16n8k16 一条 2048 MAC，wgmma m64n256k16 一条至多 262144 MAC），让取指译码等固定控制成本摊薄上千倍，并把「沿 K 维归约」直接收进指令语义；代价是操作数必须按 tile 组织——NVIDIA 让碎片跨 warp 32 线程的寄存器分布，AMX/SME 则提供显式 tile 寄存器与二维 ZA 存储——形状与布局从此写进架构契约；软件暴露沿 WMMA → mma.sync → wgmma 逐层放权，从 API 托管布局到手工编排 fence/commit/wait 异步节拍，编译器经库模板、PTX 内联汇编或 Triton 三条路径把块矩阵运算下降为单条指令；稀疏、FP8/FP4、TMA 每代微架构收益都必须经 ISA 扩口才能被软件使用，架构与 ISA 由此紧紧锁死。