SpMV 是内存带宽受限
Compulsory flop:byte 比仅 ~0.166,远低于 Machine Balance。优化重点是减少内存流量,而非减少计算。
CS267 · 第 19 课
稀疏矩阵是科学计算、图分析与机器学习的基本数据结构。SpMV 看似简单,却受限于内存带宽;理解它的存储格式、并行策略与自动调优,是掌握现代高性能计算的关键一步。
当矩阵中绝大多数元素为零(通常 >90% 甚至 >99.9%)时,我们称之为稀疏矩阵。用稠密数据结构存储它们会浪费内存与计算。稀疏矩阵并非特例——它们是现代高性能计算中最常见的矩阵类型。
Harry Markowitz 在 1950 年代构建投资组合模型时就观察到:"我们的矩阵中绝大多数系数都是零,即非零元素在矩阵中是稀疏的。"这比计算机科学正式研究稀疏矩阵算法还早了数十年。
稀疏矩阵出现的领域极为广泛:
| 领域 | 具体应用 | 稀疏性来源 |
|---|---|---|
| 气候模拟 | 大气/海洋偏微分方程离散 | 局部差分格式 |
| 图像分析 | 图像分割(Image Contour Detection) | 像素间的局部亲和矩阵 |
| Web 搜索 | Google PageRank | Web 图的邻接矩阵 |
| 推荐系统 | Netflix 评分矩阵 | 用户仅评价极少数物品 |
| 图神经网络 | GCN、GraphSAGE 的消息传播 | 图的邻接矩阵 |
| 文本分析 | 潜在语义索引(LSI) | 词-文档 bag-of-words 矩阵 |
| 结构工程 | 有限元分析(NASA 结构问题) | 局部单元连接性 |
| 线性规划 | 运筹优化 | 约束矩阵高度稀疏 |
以图像分割为例:我们希望识别图像中不同物体的边界。为此构造一个亲和矩阵(affinity matrix)\(A\),其中 \(A_{ij}\) 表示像素 \(i\) 与像素 \(j\) 属于同一物体的"亲和力"。由于只有相邻且颜色相近的像素才有关联,这个矩阵极其稀疏。最终通过求解 \(Ax = \lambda x\) 的特征问题来找到分割边界。
稀疏矩阵的非零模式(nonzero pattern)千变万化:对角、三对角、块对角、带状、完全不规则……没有一种存储格式能在所有场景下都最优。SpMV 的性能通常受限于内存带宽而非计算能力——这正是稀疏矩阵与稠密矩阵的根本区别。
稀疏矩阵存储的核心矛盾是:既要节省空间,又要支持高效的矩阵-向量乘法。不同的稀疏模式需要不同的存储策略。
CSR 是最常用的稀疏矩阵格式,由三个数组组成:
CSR 与图的邻接表有着深刻的同构关系:ptr 就是每行的邻接表起始偏移,ind 就是邻居列表,val 就是边权重。一个有向图天然对应一个稀疏矩阵,无向图对应对称矩阵(只需存储一半)。这种同构使得许多图算法可以自然地表示为稀疏线性代数运算。
| 格式 | 全称 | 特点 | 适用场景 |
|---|---|---|---|
| CSC | Compressed Sparse Column | 按列压缩,类似 CSR 的列视角 | 列操作频繁(如 MATLAB) |
| COO | Coordinate Format | (row, col, val) 三元组列表 | 构建/修改矩阵,GPU 友好 |
| DIAG | Diagonal | 按对角线存储,带偏移量 | 差分方程(有限差分/有限元) |
| Symmetric | 对称存储 | 只存上/下三角,索引复杂 | 对称矩阵(省一半空间) |
| Blocked | 分块格式 | 每个 r×c 块连续存储 | 矩阵具有天然块结构 |
CSR 本质上是把图的邻接表"压缩"进连续数组:所有行的邻居列表被拼接成一个长数组,ptr 数组记录每行在拼接数组中的起始位置。这种布局对缓存极其友好——行内访问是顺序的,且没有链表指针的开销。
SpMV(Sparse Matrix-Vector Multiplication)是稀疏线性代数中最核心的 kernel:\(y \leftarrow y + Ax\)。CSR 格式下的串行实现非常简洁:
def spmv_csr(ptr, ind, val, x, y, m):
"""y = y + A*x,A 以 CSR 格式存储,m 行"""
for i in range(m): # 遍历每一行
for k in range(ptr[i], ptr[i+1]): # 遍历该行所有非零元
y[i] += val[k] * x[ind[k]]
分析这个 kernel 的内存访问模式:
SpMV 属于 BLAS2(矩阵-向量)操作,计算强度极低:每个非零元只贡献 2 个 flop,却需要从内存加载 ~8 字节(列索引 + 值)。compulsory 的 flop:byte 比仅约 0.166,远低于典型 Machine Balance(5-10)。这意味着 SpMV 几乎总是内存带宽受限,而非计算受限。
矩阵 \(A\) 本身没有任何跨行的重用——每行读取一次就丢弃。唯一可能的重用来自 \(x\) 向量的缓存命中:如果多行访问 \(x\) 的同一位置且时间上足够接近,就能命中缓存。这就是为什么"近对角"或"良序"的矩阵 SpMV 性能更好——它们的列索引集中在较小范围内,\(x\) 的 working set 能放入缓存。
共享内存并行是最简单的 SpMV 并行化场景,因为每行的计算天然独立。
在 CSR 的串行 SpMV 上直接加 OpenMP #pragma omp for 是最自然的并行方式:
#pragma omp parallel num_threads(thread_num)
{
#pragma omp for private(j, i, tmp) schedule(static)
for (int i=0; i<m; i++) {
for (j = ptr[i]; j < ptr[i+1]; j++) {
tmp = ind[j];
y[i] += val[j] * x[tmp];
}
}
}
调度策略的选择取决于非零元的分布:
现代 CPU 提供 gather 指令(如 AVX2 的 vpgatherdd),可以从不规则地址加载向量。一个 8-wide 的 AVX2 SpMV 实现:
void avx2_csr_spmv(float *A, int32_t *nIdx, int32_t **indices,
float *x, int32_t m, float *y) {
int32_t A_offset = 0;
for(int32_t i = 0; i < m; i++) {
int32_t nElem = nIdx[i]; float t = 0.0f;
__m256 vT = _mm256_setzero_ps();
int32_t smLen = nElem - (nElem & 7); // 8 的倍数
for(int32_t j = 0; j < smLen; j+=8) {
__m256i vIdx = _mm256_loadu_si256((__m256i*)&(indices[i][j]));
__m256 vX = _mm256_i32gather_ps((float const*)x, vIdx, 4); // gather x
__m256 vA = _mm256_loadu_ps(&A[A_offset + j]); // 连续加载 A
vT = _mm256_add_ps(vT, _mm256_mul_ps(vX, vA)); // FMA
}
t += sum8(vT); // 水平归约
for(int32_t j = smLen; j < nElem; j++) // 尾部标量处理
t += x[indices[i][j]] * A[A_offset + j];
y[i] = t;
A_offset += nElem;
}
}
注意 gather 指令虽然强大,但性能仍受限于内存随机访问的延迟。当 \(x\) 的 working set 超过缓存容量时,gather 的吞吐会急剧下降。
GPU 上的朴素 CSR SpMV 让每行映射到一个线程:
// Parallel SpMV using CSR format
__global__ void spmv(int *ptr, int *ind, float *val,
int m, float *x, float *y) {
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < m; i += blockDim.x * gridDim.x) {
float yi = 0;
for (int j = ptr[i]; j < ptr[i+1]; j++) {
yi += val[j] * x[ind[j]];
}
y[i] = yi;
}
}
当行长度差异大时(如 power-law 图),朴素的一行一线程映射会导致严重的 warp 内负载不均。实践中常用warp-per-row、thread-per-row 或segmented reduction 等更精细的策略来改善。
这是一种向量机时代的经典 SpMV 算法,思想在今天仍有启发性。核心是把 CSR 的"行边界"编码为 FLAG 数组,然后用分段扫描(segmented scan)完成行内归约:
这种表达把"行边界"信息从循环控制中完全消除,转化为数据并行的扫描操作,天然适合向量处理器和 GPU 的 SIMT 执行模型。
对角格式特别适合具有带状结构的矩阵(如有限差分生成的三对角矩阵)。每条对角线存储为一个一维数组,附带一个列偏移量:
for each diagonal k do
#pragma omp parallel for
for each row i do
column = i + columnoffset[k]
if (column >= 0 && column < n)
y[i] = y[i] + val[k][column] * x[column]
这种格式的访问模式更规则(对角线上是步长为 1 的顺序访问),且天然适合 GPU 和向量机。进阶变体包括"锯齿对角"(jagged diagonal)以避免过度填充,以及按对角线长度排序以改善负载均衡。
当矩阵大到单机内存放不下时,必须跨节点分布。SpMV 的分布式并行策略与稠密矩阵既有相似之处,也有本质差异。
将矩阵 \(A\) 和结果 \(y\) 按行划分给各处理器。每个处理器 \(p\) 持有 \(A_p\) 和 \(y_p\),计算 \(y_p \leftarrow A_p x\)。关键问题:\(A_p\) 中某些列索引可能指向 \(x\) 中由其他处理器持有的分量,因此需要通信获取这些 \(x\) 的元素。
将矩阵 \(A\) 和向量 \(x\) 按列划分。每个处理器计算一个局部的 \(temp_y \leftarrow A_j x_j\),然后通过稀疏归约(sparse sum-reduce)将所有 \(temp_y\) 累加到最终的 \(y\) 上。
将处理器组织成 \(p_1 \times p_2\) 的网格。矩阵 \(A\) 也按 2D 块划分,每个处理器只与同行/同列的处理器通信。这是 NAS CG benchmark 采用的策略,在 \(p\) 很大且非零元分布均匀时效果最好。
如果矩阵能通过行列重排变成块对角形式,那么每个处理器只需持有自己的对角块和对应的 \(x, y\) 分量——完全不需要通信!图划分工具(如 METIS、Scotch)可以将稀疏矩阵的行列重新排序,使绝大多数非零元集中在对角块内,块间连接极少。
重排序不会改变数学结果,但会彻底改变通信模式。这是分布式 SpMV 性能调优中最强大的技术之一。
HPCG(High Performance Conjugate Gradients)是 TOP500 的另一个 benchmark,专门测试稀疏线性系统的求解性能。它基于一个合成的 3D 热传导 PDE(27 点 stencil,每行 27 个非零元,对称正定)。
| 系统 | HPCG [Pflop/s] | Rmax [Pflop/s] | HPCG/Peak | HPCG/HPL |
|---|---|---|---|---|
| Fugaku (A64FX) | 16 | 442 | 3% | 82% |
| Frontier (AMD EPYC) | 14 | 1353 | .7% | 66% |
| Aurora (Intel Xeon Max) | 5.6 | 1012 | .3% | — |
| LUMI (AMD EPYC) | 4.6 | 380 | .9% | — |
注意 HPCG/Typical 比值仅 0.3%-3%——这直观地说明了稀疏计算在当代超算上远远无法达到峰值。HPCG/HPL 比值(如 Fugaku 的 82%)则反映不同架构对稀疏 vs 稠密负载的适配度差异。
SpMV 的瓶颈是从内存中取矩阵数据。每个非零元只产生 2 个 flop,却需要加载列索引(4 字节)和值(4 字节)——共 8 字节。如果能减少每个非零元的索引开销,就能提升有效带宽利用率。
核心思想:不要为每个非零元存储一个列索引,而是把矩阵分成 \(r \times c\) 的小块,每个块只存一个列索引。块内的 \(r \times c\) 个元素被当作稠密小矩阵处理(无需索引)。
以一个 NASA 结构分析矩阵(n=21200, nnz=1.5M)为例:\(r=c=8\) 时存储最小,看起来是最佳选择……但真的如此吗?
在 Itanium 2 处理器上,\(r \times c = 4 \times 2\) 的块大小达到了最佳性能(31% 峰值)。更大的块(如 8×8)虽然减少了存储,但引入了大量显式零填充(fill overhead),浪费了计算和带宽。最优块大小取决于架构特性与矩阵非零模式,必须通过搜索确定。
不同架构的最优块大小差异巨大:
| 架构 | 最佳块大小 | 峰值百分比 |
|---|---|---|
| Power3 | — | 13% |
| Power4 | — | 14% |
| Itanium 2 | 4×2 | 31% |
| Itanium 1 | — | 7% |
| Ultra 2i | — | 9% |
| Pentium III | — | 19% |
大多数矩阵没有整齐的块结构。以 FEM 流体力学矩阵(N=16614, NNZ=1.1M)为例:3×3 的块看起来很自然,但许多块并不完整——为了形成完整的 r×c 块,需要填充显式零(explicit zeros)。
填充的代价:
在 Pentium III 上,1.5× 填充开销的 3×3 分块反而比无填充的实现快 1.5×——因为向量化收益超过了填充代价。这种 trade-off 强烈依赖于架构。
OSKI(以及其多核版本 pOSKI)是 Berkeley 的 BeBOP 小组开发的自动调优稀疏 kernel 库。它的核心思想是:将矩阵格式的选择与优化延迟到运行时,通过搜索找到最佳配置。
自动调优本身有成本:下界约为 5-40 次未分块 SpMV(用于填充估计),启发式策略通常只需 1-11 次。调优的收益必须超过其成本才有意义——对于会被多次重用的矩阵(如迭代求解器中的情况),调优开销可以摊销。
通过计算统计置信区间来自动控制采样比例 \(s\):监测采样方差,当方差足够小时停止采样。这避免了过度采样或采样不足。
后续研究将统计学习和贝叶斯优化引入自动调优:
对于 CSR 上的 SpMV,唯一的数据重用机会在 \(x\) 向量中。如果矩阵的列索引集中在较小范围内(如近对角矩阵),\(x\) 的 working set 能放入缓存,自然就有良好的重用。
Cache blocking 对"短宽"矩阵(行数少、列数多)有效。以 LSI(潜在语义索引)矩阵为例:
对于列索引分布极广的随机矩阵,cache blocking 几乎没有收益——因为 \(x\) 的 working set 远超缓存容量,无论怎么分块都无法改善。是否使用 cache blocking 取决于矩阵的稀疏模式。
Roofline 模型为 SpMV 提供了直观的性能分析框架:
| 优化层次 | 关键效果 |
|---|---|
| Out-of-the-box parallel | 2 个单位步长流,无 ILP/DLP,FP 占 12-25%,compulsory flop:byte < 0.166 |
| NUMA + SW prefetch | 利用所有内存通道,减少 DRAM 延迟 |
| Matrix compression (Register blocking) | 提升 ILP、DLP、flop:byte 比和 FP 占比 |
| SIMD | 每周期更多 flop |
| ILP(指令级并行) | 隐藏延迟 |
Register blocking 是 SpMV 上最有效的单优化:它同时改善了计算强度(减少索引开销)、ILP(块内循环展开)、DLP(向量化)和 FP 占比(减少整数地址计算)。实测中 register blocking 可带来最高 4× 的加速。
其他串行优化技术的典型收益:
| 优化技术 | 最高加速比 |
|---|---|
| Register blocking | 4× over CSR |
| Variable block splitting | 2.1× over CSR, 1.8× over RB |
| Diagonals | 2× over CSR |
| Reordering + splitting | 2× over CSR |
| Symmetry | 2.8× over CSR, 2.6× over RB |
| Cache blocking | 2.8× over CSR |
| Multiple vectors (SpMM) | 7× over CSR |
在实际应用中,我们很少只做一次 SpMV。迭代求解器(Iterative Solvers)的核心是一系列 SpMV 调用,用于求解 \(Ax = b\) 或 \(Ax = \lambda x\)。
Krylov 子空间方法通过构造 Krylov 子空间 \(\mathcal{K}_k(A, x) = \text{span}\{x, Ax, A^2x, \ldots, A^{k-1}x\}\),在其中寻找最佳近似解。
经典 Krylov 方法概览:
| 方法 | 适用问题 | 矩阵要求 | 核心操作 |
|---|---|---|---|
| CG (Conjugate Gradient) | \(Ax=b\) | 对称正定 (SPD) | SpMV + 向量内积 |
| GMRES | \(Ax=b\) | 一般非对称 | SpMV + Modified Gram-Schmidt |
| Lanczos | \(Ax=\lambda x\) | 对称 | SpMV + 三项递推 |
| Arnoldi | \(Ax=\lambda x\) | 一般非对称 | SpMV + MGS |
在 Krylov 方法中,\(k\) 次 SpMV 的计算占绝对主导地位,且每次 SpMV 都受限于内存带宽。当 \(k\) 很大时(如几百到几千次迭代),数据移动开销成为瓶颈。
既然 SpMV 受限于内存/通信,而 Krylov 方法要做 \(k\) 次 SpMV,能否重新组织算法,使得总通信量不依赖于 \(k\)?这就是 Communication-Avoiding (CA) 迭代求解器的核心思想。
传统 Krylov 方法每次迭代做一次 SpMV,共需 \(k\) 次通信。Matrix Powers Kernel 一次性计算 \(x, Ax, A^2x, \ldots, A^k x\) 这一整组向量,通过冗余计算换取通信减少:
关键洞察:如果矩阵"良划分"(well-partitioned,即 surface-to-volume 比小),每个处理器可以持有足够多的" ghost zone "(halo),使得连续 \(k\) 次 SpMV 中的前若干次不需要通信。
标准 GMRES 的通信成本:
CA-GMRES 用 Matrix Powers Kernel + TSQR(Tall Skinny QR)重构算法:
# 标准 GMRES(k 步)
for i in range(k):
w = A @ v[i-1] # SpMV → 1 次通信
MGS(w, v[0], ..., v[i-1]) # 正交化 → log(p) 次通信
update v[i], H
# CA-GMRES(k 步)
W = [v, Av, A^2v, ..., A^kv] # Matrix Powers Kernel → O(1) 通信
[Q, R] = TSQR(W) # Tall Skinny QR → O(log p) 通信
Build H from R, solve LSQ
通信成本对比:
| 实现 | 串行数据移动 | 并行消息数 |
|---|---|---|
| 标准 GMRES | \(O(k \cdot \nnz) + O(k^2 \cdot n)\) | \(O(k) + O(k^2 \cdot \log p)\) |
| CA-GMRES | \(O(\nnz) + O(k \cdot n)\) | \(O(1) + O(\log p)\) |
在精确算术下,CA-GMRES 与标准 GMRES 等价。但在浮点运算中,\(A^k v\) 的单项式基(monomial basis)条件数可能非常高,导致正交性丧失。实践中常用 Newton 基或 Chebyshev 基改善条件数,代价是略微增加通信。
在 8 核 Intel Clovertown 上的实测显示:CA-GMRES 在良划分矩阵上可以获得显著加速,但随着 \(k\) 增大,数值稳定性成为瓶颈。
稀疏矩阵-稀疏矩阵乘法(SpGEMM)\(C \leftarrow A \times B\) 是许多高阶算法的基础(如代数多重网格、图算法中的三角形计数、图神经网络中的邻接聚合)。与 SpMV 相比,SpGEMM 的挑战在于:
SpGEMM 的常见实现策略包括:
SpGEMM 是 GPU 上最具挑战性的稀疏 kernel 之一,也是近年来的研究热点。后续课程会更详细地讨论图算法中的 SpGEMM 应用。
Q1:CSR 格式存储一个 \(m \times n\) 的稀疏矩阵(nnz 个非零元)需要多少存储空间(32 位索引 + 32 位浮点值)?
Q2:SpMV 的 compulsory computational intensity(flop:byte ratio)约为多少?
Q3:为什么 register blocking 能提升 SpMV 性能?
Q4:CA-GMRES 相比标准 GMRES 的主要优势是什么?
Q5:分布式 SpMV 中,图划分(graph partitioning)的主要目的是:
Compulsory flop:byte 比仅 ~0.166,远低于 Machine Balance。优化重点是减少内存流量,而非减少计算。
r×c 分块可减少索引开销,最高 4× 加速。最优块大小需要通过搜索确定,依赖架构与矩阵结构。
通过图划分将近似块对角结构引入矩阵,可大幅减少处理器间通信。HPCG benchmark 揭示稀疏计算远未达峰值。
CA-GMRES 等通信避免算法通过冗余计算换通信减少,将 k 次 SpMV 的通信从 O(k) 降到 O(1)——渐近最优。