矩阵加速单元 — Tensor Core 与 XMX
更新于 2026-04-03
在 GPU Architecture 文章中，我们了解了 SM 内部的 Processing Block 结构 — 每个 Block 有 Warp Scheduler、FP32 CUDA Core、INT32 Core 和一个 Tensor Core。
那篇文章中 Tensor Core 只是一个名字。本文要回答：它内部到底是什么，为什么能让矩阵乘法快一个数量级。同时，我们会深入 Intel 的对应技术 — XMX (Xe Matrix eXtensions)，理解两大阵营如何用不同的硬件实现相同的目标。
Section 1: 为什么需要专用矩阵单元
传统 CUDA Core 做矩阵乘法是逐元素标量乘加。以一个 4×4 矩阵乘为例：
- C(4×4) = A(4×4) × B(4×4)
- 每个输出元素需要 4 次乘法 + 3 次加法 = 7 次标量操作
- 16 个输出元素 × 7 = 112 次标量操作
而 Tensor Core 一条 MMA (Matrix Multiply-Accumulate) 指令就完成整块矩阵乘累加 D = A·B + C。
吞吐量差距是数量级的：H100 SXM 上 FP32 CUDA Core 约 67 TFLOPS，FP16 Tensor Core 约 990 TFLOPS — 约 15 倍差距。
AI 训练和推理中 90%+ 的计算量是矩阵乘法（QKV 投影、Attention score、FFN 全是 GEMM）。这就是为什么专用矩阵加速单元如此重要 — 它直接决定了 AI 工作负载的性能天花板。
那么 Tensor Core 和 XMX 内部是怎么做到一拍完成矩阵乘的？答案是 Systolic Array（脉动阵列）。
Section 2: Systolic Array — 数据在阵列中脉动流动
基本概念
Systolic Array 是一种由大量简单 PE (Processing Element) 组成的计算网格。核心思想：
- 数据从边缘流入，沿固定方向在 PE 之间传递
- 每个 PE 执行一次 Multiply-Accumulate (MAC)，然后将数据传给邻居
- 数据被多个 PE 复用 — 一个输入元素经过多个 PE，参与多次计算
- 不需要每次都从内存读取 — 极高的数据复用率是效率的关键
这个名字来自 “systole”（心脏收缩），形容数据像血液一样有节奏地在阵列中脉动流动。
4×4 脉动阵列动画
下面的动画展示一个 4×4 output-stationary systolic array 的工作过程。A 矩阵的行从左侧流入，B 矩阵的列从顶部流入，每个 PE 累加自己负责的输出元素。
关键观察：
- 波前对角线 — 活跃 PE 像波浪一样从左上角扫到右下角
- PE(i,j) 在 cycle i+j+k 时处理第 k 对输入 — 输入的错开保证数据同时到达
- 4×4 矩阵乘需要 10 个 cycle（而非 1 个），但硬件可以流水线化多组矩阵
为什么不直接用 CUDA Core 并行？
看到”10 个 cycle”你可能会问：如果开 16 个 CUDA Core，每个算 C 的一个元素（4 次 MAC），不是只要 4 个 cycle 就完成了？延迟更低，为什么还需要 systolic array？
Systolic array 的优势不在延迟，而在内存访问量和面积能效：
内存访问量对比（4×4 矩阵乘）：
| 方案 | 计算方式 | 内存读取次数 |
|---|---|---|
| 16 个 CUDA Core 并行 | 每个线程读 A 的一行（4 个）+ B 的一列（4 个） | 16 × 8 = 128 次 |
| 4×4 Systolic Array | A 的每个元素从左侧流入，流过 4 个 PE 被复用；B 同理 | 16 + 16 = 32 次 |
A 和 B 总共只有 32 个不同的元素，但 16 个 CUDA Core 各自独立读取，产生了 4 倍的冗余访问。
你可能会想到：用 shared memory 缓存 A 和 B 的 tile，让线程从 shared memory 读取而不是从 global memory 重复读。这确实是 GEMM 优化的标准做法，但它引入了额外开销：
// 阶段 1：所有线程协作加载 tile 到 shared memory
shared_A[ty][tx] = A[row][tx];
shared_B[ty][tx] = B[ty][col];
__syncthreads(); // ← barrier 同步
// 阶段 2：从 shared memory 读取计算
sum += shared_A[ty][k] * shared_B[k][tx];__syncthreads() 是一个 barrier（屏障） — thread block 内的所有线程必须都执行到这一行才能继续。为什么需要它？因为实际 GEMM 中一个 thread block 远不止 16 个线程 — 典型配置是 256 个线程，分布在 8 个 warp 上。同一 warp 内的 32 个线程锁步执行，但不同 warp 之间由 warp scheduler 独立调度，执行进度完全可能不同。Warp 0 的线程可能需要读取 Warp 3 写入 shared memory 的数据，而 Warp 3 可能还没执行到写入那一行。barrier 确保所有 warp 都完成写入后才开始读取。
即使用了 shared memory + barrier，所有线程仍然需要大量 shared memory 读取（只是从 global memory 变成了 shared memory），加上 barrier 等待的同步开销。
Systolic array 完全不需要这些 — 数据按硬件固定的时序在相邻 PE 间传递，复用是物理上天然发生的，没有同步指令、没有 shared memory 访问。
关于 thread block 与 warp 的关系、shared memory 和寄存器的作用域、
__syncthreads()的工作原理等编程模型细节，参见 GPU Architecture — 硬件概念 vs 软件抽象 和 CUDA 编程模型。
面积与能效：
| CUDA Core | Systolic Array PE | |
|---|---|---|
| 组成 | 取指、译码、寄存器堆、ALU、分支… | 一个 MAC 单元 + 寄存器 |
| 能力 | 通用：可以执行任意指令 | 专用：只做乘累加 |
| 单元面积 | 大 | 小（同面积可放更多） |
| 数据搬运能耗 | shared memory 访问 ~5 pJ | 相邻 PE 寄存器传递 ~0.5 pJ |
同样的芯片面积下，systolic array 可以塞进更多 MAC 单元，并用更少的带宽和能耗喂饱它们。这就是为什么 Tensor Core 吞吐量能达到 CUDA Core 的约 15 倍 — 不是因为单次操作更快，而是因为专用硬件在面积和带宽上都更高效。
4×4 看起来差距不大，但放到实际尺寸（如 4096×4096 的 GEMM）：naive 并行方案的冗余内存访问量与矩阵维度成正比增长，带宽瓶颈会成为决定性限制。
Dataflow 变体
Systolic array 有多种数据流模式，区别在于哪个矩阵”固定”在 PE 内，哪个”流动”：
- Output-Stationary: 输出矩阵 C 固定在 PE 中累加，A 和 B 都流过。优点是部分和不需要移动
- Weight-Stationary: 权重矩阵 B 预加载到 PE，输入 A 流过，部分和向下传递。适合推理（权重固定）
Tensor Core 和 XMX 的具体实现是厂商机密，但本质都是 systolic array 的变体。
Section 3: NVIDIA Tensor Core
MMA 操作
Tensor Core 的核心操作是矩阵乘累加：D = A × B + C
- 原始概念尺寸是 4×4，但从 Volta 到 Blackwell，实际支持的块尺寸越来越大
- 尺寸标注遵循 BLAS GEMM 约定：，其中 m 是输出行数，n 是输出列数，k 是内积维度（A 的列数 / B 的行数）。例如
m16n16k16表示 A(16×16) × B(16×16) → D(16×16)，m16n8k16表示 A(16×16) × B(16×8) → D(16×8) - Hopper (4th gen) wmma API 常用尺寸为
m16n16k16；PTX 层面有更小的m16n8k16 - 这些是 Tensor Core 单条指令能处理的矩阵块大小。更大的矩阵乘需要软件做 tiling，将大矩阵拆成这些小块反复调用
- 每个 SM 有 4 个 Tensor Core（每个 Processing Block 一个）
精度支持演进
从 Volta (2017) 到 Blackwell (2024)，Tensor Core 支持的精度不断扩展：
趋势非常清晰：精度越来越低，吞吐量越来越高。FP8 的吞吐量是 FP16 的 2 倍，FP4 又翻倍。这也是为什么 FP8 训练（如 DeepSeek V3）和 FP4 量化推理成为趋势。
注意：Turing/Ampere 时代的 INT4 和 INT1 支持在 Hopper 之后被移除 — 实际 AI 工作负载更青睐 FP8/FP4 等浮点低精度格式。
Warp 级操作
Tensor Core 操作不是单个线程发起的 — 它是 warp 级协作操作。一个 warp 的 32 个线程共同持有输入矩阵的片段（fragment），一起发射 MMA 指令。
关键要点：
- Fragment 是矩阵块在 32 个线程寄存器中的分布式表示 — 每个线程只持有一部分
wmma::load_matrix_sync从内存加载到 fragment，wmma::mma_sync执行矩阵乘，wmma::store_matrix_sync写回内存- wmma API 操作 m16n16k16 块；底层 PTX 指令
mma.sync.aligned.m16n8k16操作更小的块，由编译器拆分
Section 4: Intel XMX
Intel 的矩阵加速单元叫 XMX (Xe Matrix eXtensions)，是 Xe2 架构的核心组件。
关键规格 (Xe2 / Lunar Lake)
- 每个 Xe-Core 包含 8 个 XMX 单元 + 2 个 Vector Engine
- XMX 内部是 8×8 systolic array
- 支持精度：FP16、BF16、TF32、INT8、INT4
- 编程接口：SYCL
joint_matrix(高层) / ESIMDdpas(低层)
dp4a vs dpas — 容易混淆的两条指令：
- dp4a (Dot Product of 4 elements and Accumulate)：计算 4 对 INT8 元素的点积并累加到 INT32（
acc += a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3]）。在 Vector Engine 上执行，不走 XMX。本质是一条向量点积指令，类似 NVIDIA 的__dp4aintrinsic（Pascal/Turing 开始支持）。dp4a 在 Xe 架构之前就存在- dpas (Dot Product Accumulate Systolic)：驱动 XMX 脉动阵列的矩阵指令，一条指令完成一个小矩阵块的乘累加，类似 NVIDIA 的
mma.sync（Tensor Core 指令）。吞吐量远高于 dp4a，因为整个 8×8 systolic array 同时工作简言之：dp4a 是标量点积（Vector Engine），dpas 是矩阵乘累加（XMX systolic array）。
与 NVIDIA 的架构类比
| NVIDIA | Intel | 说明 |
|---|---|---|
| SM | Xe-Core | 基本计算单元 |
| CUDA Core | Vector Engine | 标量/向量计算 |
| Tensor Core | XMX | 矩阵加速 |
| Warp (32 threads) | Sub-group (8/16 wide) | SIMD 执行单位 |
| Shared Memory | SLM (Shared Local Memory) | SM/Xe-Core 内共享存储 |
| wmma / mma.sync | joint_matrix / dpas | 矩阵操作 API |
核心思路完全一致 — 都是用 systolic array 做 D = A × B + C。主要区别在规模（数据中心 GPU vs 客户端 iGPU）和编程模型（CUDA 的 warp-level vs SYCL 的 sub-group）。
Section 5: Tensor Core vs XMX 详细对比
| 对比维度 | NVIDIA Tensor Core H100 Hopper | Intel XMX Xe2 Lunar Lake |
|---|---|---|
| 厂商 / 架构 | NVIDIA (Hopper, 4th gen) | Intel (Xe2, Lunar Lake) |
| 内部结构 | Systolic Array 变体 | Systolic Array (8×8)(相似) |
| 核心操作 | D = A × B + C (MMA) | D = A × B + C (DPAS)(相似) |
| 矩阵块尺寸 (FP16) | m16n8k16 | m8n8k16 |
| 每 SM/Xe-Core 数量 | 4 Tensor Core / SM | 8 XMX / Xe-Core |
| FP16 / BF16 | 990 TFLOPS (H100 SXM) | ~48 TOPS (Lunar Lake iGPU) |
| FP8 支持 | Hopper 起 (4th gen) | Xe2 起 (Lunar Lake) |
| FP4 支持 | Blackwell 起 (5th gen) | 尚未支持 |
| TF32 支持 | Ampere 起 (3rd gen) | Xe2 起 |
| 编程接口 (高层) | CUDA wmma / mma.sync | SYCL joint_matrix |
| 编程接口 (低层) | PTX mma 指令 | ESIMD dpas 指令 |
| Warp/Sub-group 协作 | 32 线程 warp 协作 | 8/16-wide sub-group |
| 目标场景 | 数据中心 AI 训练/推理 | 客户端 AI 推理 (iGPU) |
两者核心思路相同: systolic array 做 D=A×B+C。主要区别在规模 (数据中心 vs 客户端)、矩阵块尺寸和编程接口。Hover 高亮行。
两者的相似点比区别更重要：
- 都是 systolic array 变体，都做 D = A·B + C
- 都只能做特定尺寸、特定精度的矩阵乘 — 其他操作仍走传统 CUDA Core / Vector Engine
- 都需要软件配合（tiling、数据对齐）才能发挥峰值性能 — 这是 GEMM 优化文章的主题
关键区别主要在规模和目标场景：
- H100 Tensor Core 峰值 ~990 TFLOPS (FP16)，面向数据中心 AI 训练
- Lunar Lake XMX 峰值 ~48 TOPS (INT8)，面向客户端 AI 推理
- NVIDIA 的 32-wide warp vs Intel 的 8/16-wide sub-group 影响编程方式
Section 6: 跨厂商标准 — Subgroup 与 Cooperative Matrix
前面介绍的 wmma（NVIDIA）和 joint_matrix（Intel SYCL）是厂商专有的 API。Vulkan / SPIR-V / OpenCL 等跨平台标准则定义了统一的抽象，使同一套代码可以运行在不同厂商的硬件上。
Subgroup：warp / wave 的跨厂商统一术语
Subgroup 是指一组在硬件上以锁步方式（SIMD/SIMT）执行的线程。这个概念在不同厂商有不同的叫法：
| 厂商 | 术语 | 典型宽度 |
|---|---|---|
| NVIDIA | Warp | 32 线程 |
| AMD | Wave / Wavefront | 32 或 64 线程 |
| Intel | Sub-group | 8 / 16 / 32 线程 |
| Vulkan / SPIR-V | Subgroup | 取决于硬件 |
Subgroup 内的线程可以直接交换寄存器数据，不需要经过 shared memory，支持的操作包括：
- Shuffle：线程 A 直接读线程 B 的寄存器值
- Broadcast：一个线程的值广播给 subgroup 内所有线程
- Reduction：subgroup 内直接求和 / 求最大值
这比走 shared memory（store → barrier → load）快得多，是高性能 GPU kernel 的基础原语。
Cooperative Matrix：Tensor Core / XMX 的跨厂商编程抽象
前文我们看到，Tensor Core 的 MMA 操作是 warp 级协作 — 32 个线程共同持有矩阵片段（fragment），一起发射矩阵乘指令。Cooperative matrix 就是将这种”subgroup 内多线程合作持有并计算矩阵”的模式标准化：
- 没有任何单个线程持有完整矩阵，矩阵的片段分散在 subgroup 内各线程的寄存器中
- 硬件以 subgroup 为粒度执行矩阵乘累加
- 映射到底层硬件：NVIDIA Tensor Core、Intel XMX、AMD Matrix Core
各层 API 的对应关系：
| 抽象层级 | NVIDIA | Intel | 跨厂商标准 |
|---|---|---|---|
| 高层 API | wmma::mma_sync | joint_matrix_mad | cooperativeMatrixMulAdd (Vulkan) |
| 加载 | wmma::load_matrix_sync | joint_matrix_load | cooperativeMatrixLoad |
| 存储 | wmma::store_matrix_sync | joint_matrix_store | cooperativeMatrixStore |
| 规范 | CUDA wmma API | SYCL joint_matrix | VK_KHR_cooperative_matrix / SPV_KHR_cooperative_matrix |
典型的使用流程（以 Vulkan cooperative matrix 为例）：
1. cooperativeMatrixLoad — 从显存/shared memory 加载矩阵片段到 subgroup 的寄存器中
2. cooperativeMatrixMulAdd — 硬件执行 D = A × B + C（一条指令完成小矩阵乘加）
3. cooperativeMatrixStore — 将结果写回显存/shared memory矩阵的可用尺寸取决于硬件和数据类型。程序可以在运行时查询硬件支持的尺寸组合（如 NVIDIA FP16 常见 16×16×16，Intel 可能支持 8×16×32 等）。
为什么需要跨厂商标准
厂商专有 API（wmma、joint_matrix）通常更成熟、优化更深入。但 Vulkan cooperative matrix 的价值在于：
- 可移植性：同一份 SPIR-V shader 可以在 NVIDIA、AMD、Intel 的 GPU 上运行
- 生态统一：推理框架（如使用 Vulkan 后端的 llama.cpp / GGML）不需要为每个厂商维护独立的矩阵乘 kernel
- 覆盖面广：Vulkan 支持桌面、移动、嵌入式等各种平台
实际性能上，跨厂商标准的 overhead 通常很小 — 底层驱动会将 cooperative matrix 操作映射到硬件原生的矩阵指令（Tensor Core HMMA、Intel DPAS 等）。
Section 7: Dual-Pipe — 同时利用两种计算单元
Tensor Core 和 CUDA Core 是 SM 内不同的功能单元，有独立的执行通道。传统做法是串行执行：矩阵乘（Tensor Core）完成后才开始 element-wise 操作（CUDA Core）。
Dual-Pipe 优化打破了这个限制：将输入拆成多个 micro-batch，当 Tensor Core 处理batch B 的 GEMM 时，CUDA Core 同时处理batch A 的 activation/normalization。不同 batch 之间没有数据依赖，因此可以安全重叠。
Dual-pipe 的条件：
- 不同 micro-batch 之间没有数据依赖 — Tensor Core 和 CUDA Core 处理的是不同 batch 的数据
- 需要精心安排 micro-batch 的调度顺序和数据布局
- 可以显著提升 SM 利用率，特别是在 MoE 等 element-wise 操作较多的架构中
总结
矩阵加速单元的核心设计思想：
- 专用硬件做矩阵乘 — Systolic array 用极高的数据复用率实现单条指令完成整块矩阵乘累加
- 精度换吞吐 — 从 FP16 到 FP8 到 FP4，每降一级精度吞吐量翻倍，AI 工作负载可以接受较低精度
- 协作执行 — 矩阵操作是 warp/sub-group 级别的协作，需要软件配合（fragment 管理、tiling 策略）
下一篇文章将从编程者的视角出发 — CUDA 编程模型，理解 thread/block/grid 的层级、shared memory 的使用、memory coalescing 等关键概念。这些知识是写高性能 GPU 代码（包括利用 Tensor Core）的前提。