第二十五课：稀疏计算、不规则访存，以及 GPU 不擅长什么

导语：这一课是倒数第三课，重点解释“计算量更少”为什么不一定更快，以及不规则访存、原子冲突和负载不均衡为何会削弱 GPU 优势。

稠密矩阵中的大部分元素都非零，例如：

A= \begin{bmatrix} 1&2&3\\ 4&5&6\\ 7&8&9 \end{bmatrix}

稀疏矩阵则包含大量零元素：

A= \begin{bmatrix} 1&0&0&0\\ 0&0&2&0\\ 0&3&0&0\\ 0&0&0&4 \end{bmatrix}

直觉上，既然零元素不需要参与乘法，稀疏计算应该比稠密计算快很多。

但在 GPU 上：

跳过零元素能够减少算术量，却可能引入索引、随机访存、分支、负载不均衡和原子冲突。

所以“稀疏”并不自动意味着“高性能”。

一、什么是稀疏度

假设矩阵共有 $N$ 个元素，其中非零元素数量为：

NNZ

稀疏度可以定义为：

\text{Sparsity} = 1-\frac{NNZ}{N}

例如一个矩阵有100万个元素，其中只有10万个非零元素：

\text{Sparsity} = 1-\frac{100000}{1000000} =90\%

对应密度为：

\text{Density}=10\%

需要注意，90%稀疏并不意味着一定能达到10倍加速。

二、为什么稠密矩阵计算很适合 GPU

稠密矩阵乘法具有非常规则的数据结构：

矩阵尺寸固定
元素连续存储
每个 Tile 工作量接近
访存地址可预测
线程控制流一致
数据可以反复复用
Tensor Core 可以高效工作

例如 GEMM：

C=AB

可以规则地分块：

A Tile × B Tile
       ↓
固定数量 MMA
       ↓
C Tile

每个线程块的工作量相近，GPU 很容易保持大量执行单元忙碌。

三、为什么稀疏计算更困难

稀疏矩阵只保存非零元素。

除了值本身，还必须保存：

非零元素位于哪一行；
位于哪一列；
每行有多少元素；
下一行从哪里开始。

因此需要额外的索引数据。

执行过程可能变成：

读取索引
  ↓
根据索引计算地址
  ↓
随机读取对应数据
  ↓
执行少量乘加

和稠密计算相比，算术指令减少了，但控制和访存开销增加了。

四、COO 格式

COO 是 Coordinate Format，坐标格式。

每个非零元素保存：

row index
column index
value

例如矩阵：

A= \begin{bmatrix} 5&0&0\\ 0&0&8\\ 0&3&0 \end{bmatrix}

COO 可以表示为：

row = [0, 1, 2]
col = [0, 2, 1]
val = [5, 8, 3]

每个非零元素是一个三元组：

(0,0,5)
(1,2,8)
(2,1,3)

COO 的优点

表示简单；
容易构造；
适合输入和增量插入；
对无序稀疏数据友好；
易于并行遍历所有非零元素。

COO 的缺点

每个元素都要保存行、列两个索引；
索引开销较大；
同一行元素可能分散；
计算输出时常需要原子累加；
局部性通常较差。

五、CSR 格式

CSR 是 Compressed Sparse Row，压缩稀疏行格式。

它使用三个数组：

values
column_indices
row_offsets

仍以矩阵：

A= \begin{bmatrix} 5&0&0\\ 0&0&8\\ 0&3&0 \end{bmatrix}

为例：

values         = [5, 8, 3]
column_indices = [0, 2, 1]
row_offsets    = [0, 1, 2, 3]

row_offsets[i] 表示第 $i$ 行非零元素在 values 中的起始位置。

例如第1行：

start = row_offsets[1] = 1
end   = row_offsets[2] = 2

所以只包含：

values[1] = 8
col[1]    = 2

CSR 的优点

按行访问方便；
同一行非零元素连续存储；
行指针开销较低；
适合矩阵乘向量；
适合逐行归约。

CSR 的缺点

不同行的非零数量可能差别很大；
列访问不方便；
仍然需要列索引；
对动态插入不友好；
行长度不均匀会造成负载不平衡。

六、CSC 格式

CSC 是 Compressed Sparse Column，压缩稀疏列格式。

它与 CSR 类似，只是按列组织：

values
row_indices
column_offsets

CSR 适合按行计算，CSC 适合按列计算。

例如：

CSR：
容易计算 A × x

CSC：
容易按列读取 A

选择哪种格式取决于算法的数据流。

七、稀疏矩阵向量乘 SpMV

SpMV 是：

y=Ax

其中：

$A$ 是稀疏矩阵；
$x$ 是稠密向量；
$y$ 是输出向量。

使用 CSR 时，第 $i$ 行为：

y_i= \sum_{j=row\_ptr[i]}^{row\_ptr[i+1]-1} values_j\cdot x_{col_j}

伪代码：

for each row i:
    sum = 0

    for p = row_ptr[i] to row_ptr[i + 1]:
        sum += values[p] * x[col_idx[p]]

    y[i] = sum

八、SpMV 的算术强度为什么很低

每个非零元素通常需要：

读取一个 value；
读取一个列索引；
根据列索引读取一个 $x$ ；
执行一次乘加。

以 FP32 和32位索引估算：

value：4 Bytes
column index：4 Bytes
x[col]：约4 Bytes

总计至少约：

12\text{ Bytes}

执行：

2\text{ FLOPs}

算术强度约：

AI\approx\frac{2}{12} \approx0.167\text{ FLOPs/Byte}

还未计入：

行指针；
输出；
Cache Miss；
无效事务。

因此 SpMV 通常是明显的内存受限算子。

九、SpMV 中最麻烦的访问

values[p] 和 col_idx[p] 在 CSR 中通常连续。

但：

x[col_idx[p]]

是间接访问。

列索引可能是：

1, 1000, 7, 500000, 23...

于是 Warp 中不同线程可能访问完全不同的地址。

结果是：

内存事务难以合并；
Cache 局部性差；
TLB 压力增加；
Long Scoreboard 增加；
显存带宽难以充分利用。

这类模式称为：

Gather，收集式访问。

十、Gather 是什么

Gather 指多个线程根据索引，从不同地址读取数据。

例如：

value = x[index[i]];

如果：

index = [0,1,2,3,4,...]

访问连续，效率较高。

如果：

index = [100,7,80000,3,...]

访问随机，效率较低。

Gather 的核心问题是：

线程连续
不代表地址连续

所以即使线程编号规则，内存访问仍然可能无法合并。

十一、Scatter 是什么

Scatter 指线程根据索引，把结果写入不同地址：

output[index[i]] += value;

如果多个线程写入不同位置，问题较小。

但如果多个线程写入同一个位置：

Thread 0 → output[5]
Thread 1 → output[5]
Thread 2 → output[5]

就会发生写冲突。

此时通常需要：

atomicAdd(&output[index[i]], value);

所以 Scatter 常伴随原子操作。

十二、为什么原子冲突会降低性能

原子操作本身未必很慢。

真正的问题是大量线程竞争同一地址。

例如图计算中，许多源顶点可能同时更新同一个目标顶点：

Vertex A ─┐
Vertex B ─┼→ Vertex X
Vertex C ─┤
Vertex D ─┘

对应：

atomicAdd(&feature[X], contribution);

这些更新必须以受控方式完成，地址热点会导致请求排队。

如果更新地址分散：

Thread 0 → output[0]
Thread 1 → output[1]
Thread 2 → output[2]

原子吞吐通常会好得多。

十三、SpMV 如何映射线程

最简单的方式是：

一个线程处理一行

Kernel 概念上为：

int row = blockIdx.x * blockDim.x
        + threadIdx.x;

float sum = 0.0f;

for (int p = row_ptr[row];
     p < row_ptr[row + 1];
     ++p) {

    sum += values[p] * x[col_idx[p]];
}

y[row] = sum;

这种方法简单，但负载不均衡非常严重。

十四、为什么“一线程一行”会负载不均衡

假设一个 Warp 中32个线程分别处理32行。

每行非零元素数量为：

Row 0：3
Row 1：2
Row 2：2000
Row 3：1
...

处理短行的线程很快完成，但 Warp 必须等待处理2000个元素的线程。

示意：

Thread 0：██
Thread 1：█
Thread 2：████████████████████
Thread 3：█

其他 Lane 虽已无工作，Warp 仍未结束。

这会造成严重的 SIMT 利用率下降。

十五、一 Warp 处理一行

对于较长行，可以让一个 Warp 协同处理一行。

例如：

Lane 0 处理第 0、32、64... 个非零元素
Lane 1 处理第 1、33、65... 个非零元素
...

每个 Lane 得到局部和，再用 Warp Shuffle 归约：

32个局部和
    ↓
Warp Reduction
    ↓
一行输出

优点：

长行内部并行；
同一 Warp 工作更一致；
归约可以用 Shuffle 完成。

缺点：

如果一行只有2个非零元素，30个 Lane 空闲；
行长差异仍然存在；
x[col] 访问仍可能随机。

十六、一 Block 处理一行

对于非常长的行，可以让整个 Block 处理一行：

多个 Warp
共同遍历一行
    ↓
Block Reduction

适合：

超长行；
图中的高度节点；
少数特别密集的行。

但如果用于所有行，短行会浪费大量线程。

因此高性能 SpMV 常根据行长度采用不同策略。

十七、自适应行分配

可以把行按长度分类：

短行：一线程一行
中等行：一 Warp 一行
长行：一 Block 一行
超长行：拆成多个任务

这种方法提高负载平衡，但增加了：

预处理；
行分类；
多个 Kernel；
任务队列；
实现复杂度。

因此稀疏计算经常需要根据具体矩阵结构定制。

十八、将非零元素均匀分配给线程

另一种思路不是按行分配，而是把全部 $NNZ$ 均匀分给线程。

例如：

Thread 0：处理非零元素 0～7
Thread 1：处理非零元素 8～15
...

这样每个线程工作量更均匀。

但一个线程的范围可能跨越多行，多个线程也可能同时处理同一行。

最终合并输出时可能需要：

分段归约；
Prefix Sum；
原子加法；
Carry 处理。

负载更均匀，但结果管理更复杂。

十九、什么是 SpMM

SpMM 是稀疏矩阵乘稠密矩阵：

C=AB

其中：

A\in\mathbb{R}^{M\times K}

是稀疏矩阵，

B\in\mathbb{R}^{K\times N}

是稠密矩阵。

与 SpMV 相比，SpMM 对每个非零元素要处理一整段 $B$ 的行：

A_{ik}\times B_{k,:}

因此一个非零权重可以复用于多个输出通道。

二十、SpMM 为什么比 SpMV 更适合 GPU

SpMV 中一个非零元素只执行一次乘加。

SpMM 中一个非零元素可对应：

N

次乘加。

因此算术强度更高。

例如：

读取一个稀疏 value
读取一个 column index
      ↓
与 B 中一整行多个元素相乘

索引开销可以被多次计算摊薄。

当 $N$ 足够大时，SpMM 通常比 SpMV 更容易利用 GPU。

二十一、但 SpMM 仍然存在的问题

虽然计算量增加，但仍有：

稀疏索引读取；
不规则的 $B$ 行访问；
行长度差异；
输出累加冲突；
稀疏 Tile 填充不足；
Tensor Core 难以直接使用。

因此性能通常仍低于尺寸相近的稠密 GEMM 峰值。

二十二、非结构化稀疏

非结构化稀疏表示零元素可以出现在任意位置：

优点：

可以获得很高稀疏率；
剪枝自由度高；
更容易保留模型精度。

缺点：

索引复杂；
访存不规则；
难以映射到 Tensor Core；
负载不均衡；
硬件利用率低。

二十三、结构化稀疏

结构化稀疏要求零元素遵守某种固定规则。

例如：

整行剪枝；
整列剪枝；
通道剪枝；
Block 稀疏；
固定 N:M 稀疏。

例如 Block Sparse：

非零 Block   零 Block
零 Block     非零 Block

Block 内部仍然是稠密的。

这样可以：

按块执行 GEMM；
减少索引数量；
提高连续访问；
利用矩阵计算单元。

二十四、2:4 结构化稀疏

2:4 稀疏表示每连续4个元素中，恰好有2个非零元素。

例如：

[3, 0, 7, 0]
[0, 5, 0, 2]

不是任意50%稀疏，而是局部固定模式：

每组4个
保留2个

硬件可以：

保存两个非零值；
保存较少的元数据；
跳过另外两个乘法；
保持规则的数据流。

二十五、为什么结构化稀疏更容易加速

硬件最怕的是不可预测。

结构化稀疏提供了确定规则：

固定分组
固定非零数量
固定计算节拍
固定输出位置

因此可以构造流水线：

读取压缩值和少量元数据
      ↓
选择对应输入
      ↓
执行固定数量乘法

执行单元不会因为某一行有3个元素、另一行有2000个元素而严重失衡。

二十六、为什么50%稀疏不一定获得2倍端到端加速

即使乘法数量减少一半，仍然可能需要：

读取稀疏元数据；
解码非零位置；
读取激活；
写输出；
执行 Bias、Activation、Norm；
完成通信；
处理不支持稀疏的其他算子。

所以端到端时间为：

T= T_{\text{sparse GEMM}} + T_{\text{other}}

只有稀疏 GEMM 部分可能加速。

根据 Amdahl 定律，如果 GEMM 只占总时间的60%，即使 GEMM 加速2倍：

\text{Speedup} = \frac{1}{0.4+0.6/2} \approx1.43

总体也只有约1.43倍。

二十七、参数减少也不等于访存按比例减少

稀疏格式需要额外存储：

非零值
索引
元数据
行指针
Padding

假设 FP16 权重2字节，32位列索引4字节。

非结构化 CSR 中，每个非零元素可能需要：

2 Bytes value
+
4 Bytes index

索引甚至比权重本身更大。

对于低精度权重，例如4位权重，索引开销会更加突出。

因此低位宽和非结构化稀疏有时并不容易同时获得良好压缩效率。

二十八、Block Sparse 的权衡

Block Sparse 将矩阵划分成固定大小的块，例如：

16\times16

某个 Block 要么全部保留，要么全部为零。

优点：

每个非零 Block 内可使用稠密计算；
索引开销被整个 Block 摊薄；
访存更连续；
便于 Tensor Core 映射。

缺点：

剪枝粒度较粗；
可能保留许多本来可以删除的元素；
相同稀疏率下模型精度可能更难保持；
Block 尺寸需要与硬件匹配。

二十九、稀疏卷积

稀疏卷积可能指不同概念。

权重稀疏

卷积核中大量权重为0。

激活稀疏

ReLU 等操作使大量激活为0。

空间稀疏

例如点云和体素中，只有少量空间位置存在有效数据。

这三类稀疏的数据结构和加速方式完全不同。

三十、激活稀疏为什么更难预测

权重在推理期间通常固定，可以预先压缩和排序。

激活由输入决定：

不同样本
不同层
不同 Token

稀疏位置随运行时变化。

因此需要动态执行：

检测零元素；
压缩；
生成索引；
调度非零任务；
解压或 Scatter 输出。

如果检测和压缩开销过大，跳零计算的收益可能被抵消。

三十一、事件驱动计算

在 SNN 等模型中，Spike 为0时可以跳过对应计算。

例如：

y+=s\cdot w

当：

s=0

时无需执行乘加。

事件驱动架构可以只处理：

s = 1

的位置。

但同样会出现：

Spike 分布不均；
事件队列；
随机权重读取；
不同神经元工作量不同；
累加地址冲突。

所以稀疏 SNN 是否高效，取决于是否有适合稀疏事件流的硬件数据通路，而不是只看脉冲稀疏率。

三十二、图计算为什么不规则

图由顶点和边组成：

G=(V,E)

不同顶点的度数可能差异很大。

例如社交网络：

普通用户：几十条边
热门账号：数百万条边

算法可能需要遍历：

一个顶点的所有邻居

因此工作量高度不均匀。

图中邻居编号通常也不连续，导致随机访存。

三十三、图的 CSR 表示

图的邻接表可以使用类似 CSR 的结构：

neighbors
vertex_offsets

例如：

Vertex 0 的邻居位于：
neighbors[offset[0] : offset[1]]

Vertex 1 的邻居位于：
neighbors[offset[1] : offset[2]]

这种表示紧凑，但不同顶点的邻居数量可能差异巨大。

所以图计算与 SpMV 的问题非常相似：

行长度不均；
Gather；
Scatter；
原子操作；
不规则控制流。

三十四、Push 与 Pull

图算法常有两种更新方向。

Push

活跃源顶点向邻居发送更新：

Source
  ├→ Neighbor A
  ├→ Neighbor B
  └→ Neighbor C

优点：

只遍历活跃源顶点的边。

缺点：

多个源可能写同一目标；
需要 Atomic；
Scatter 不规则。

Pull

每个目标顶点从入邻居收集信息：

Target
  ← Source A
  ← Source B
  ← Source C

优点：

每个线程只写自己的目标；
减少输出原子冲突。

缺点：

可能检查大量非活跃邻居；
Gather 不规则；
输入访问随机。

不同迭代阶段可能适合不同模式。

三十五、Frontier 是什么

在 BFS 等图算法中，并不是所有顶点每轮都活跃。

当前需要处理的顶点集合称为 Frontier。

例如：

Level 0：{A}
Level 1：{B,C}
Level 2：{D,E,F,G}

每一轮 Frontier 大小可能剧烈变化。

当 Frontier 很小时：

GPU 并行度不足；
Kernel 启动开销明显。

当 Frontier 很大时：

边遍历和原子冲突增加。

这使图算法性能在不同迭代间变化很大。

三十六、动态任务队列

为处理不规则任务，可以使用 GPU 工作队列：

全局任务队列
      ↓
线程块不断领取任务
      ↓
处理完成后继续领取

这种 Persistent Kernel 可以改善负载平衡。

但需要：

原子队列指针；
任务分配；
队列容量管理；
终止检测；
跨 Block 协调。

任务过小时，队列管理成本可能超过计算成本。

三十七、负载均衡的几个层次

稀疏任务可以在不同层次均衡：

静态均衡

编译或预处理时根据行长度分配任务。

动态均衡

运行时线程块从队列领取任务。

Warp 内均衡

空闲 Lane 帮助处理同一行或邻接表。

Block 间工作窃取

完成较早的 Block 继续处理其他 Block 剩余任务。

动态性越强，均衡效果可能越好，但控制开销也越高。

三十八、分支发散在稀疏算法中的表现

例如：

if (value != 0) {
    process(value);
}

同一 Warp 中：

Lane 0：非零，执行
Lane 1：零，跳过
Lane 2：零，跳过
Lane 3：非零，执行

虽然跳过了部分计算，但执行路径可能变得稀疏。

如果只有4个 Lane 非零：

4个 Lane 工作
28个 Lane 被屏蔽

那么理论上跳过了大量运算，但 SIMD 利用率也很低。

这就是为什么在普通稠密 GPU 上直接使用：

if (x != 0)

不一定能获得与稀疏度成比例的加速。

三十九、压缩后再计算是否更好

可以先把非零元素压缩：

原始稠密数据
      ↓
检测非零
      ↓
Prefix Sum
      ↓
Compact
      ↓
稀疏计算

这样后续 Warp 中大部分线程都处理非零元素。

但前处理本身需要：

扫描所有元素；
条件标记；
Prefix Sum；
Scatter；
生成索引。

如果稀疏数据只使用一次，压缩成本可能不划算。

如果同一稀疏结构会重复使用很多次，例如固定权重，则预压缩更有价值。

四十、稀疏计算适合什么情况

稀疏化更可能带来收益的条件包括：

稀疏率足够高；
稀疏结构重复使用；
格式与硬件匹配；
索引开销较低；
非零分布较规则；
计算可以按 Block 批量执行；
中间数据不需要频繁压缩和解压；
稀疏算子占端到端时间比例较高。

四十一、稀疏计算不适合什么情况

以下情况可能不如直接做稠密计算：

稀疏率只有10%～30%；
非零位置完全随机；
权重是INT4但索引是INT32；
矩阵尺寸很小；
每次输入稀疏结构都变化；
需要频繁格式转换；
稠密 Tensor Core 吞吐率极高；
稀疏 Kernel 实现不成熟；
下游算子仍要求稠密输出。

有时即使执行了大量乘零操作，稠密 GEMM 仍然更快，因为它的数据流更规则。

四十二、GPU 最擅长的任务

GPU 最擅长的是：

大量独立或规则的并行任务
相同或相近的控制流
连续、可合并的内存访问
较高的数据复用
充足的任务规模
规则的矩阵和张量运算

典型例子：

GEMM；
卷积；
图像处理；
大规模 Elementwise；
规则 Reduction；
Attention；
物理网格计算。

四十三、GPU 不擅长的任务

1. 强串行依赖

例如：

下一步必须等待上一步结果

长依赖链无法获得大量线程并行。

2. 大量复杂分支

每个线程走完全不同的路径，容易产生 Warp 发散。

3. 随机细粒度访存

例如指针链：

node = node->next;

地址难以预取和合并。

4. 工作量差异巨大

一些线程只做一次运算，另一些线程做上万次。

5. 任务规模很小

只有几十或几百个操作时，Kernel 启动和数据搬运成本过高。

6. 高频全局同步

普通 GPU Block 之间缺乏廉价的全局屏障。

7. 大量热点原子操作

许多线程竞争同一地址，更新被序列化。

8. 延迟敏感的单线程任务

GPU 追求吞吐率，不追求单线程最低延迟。

四十四、CPU 为什么可能更适合不规则任务

CPU 具有：

强大的分支预测；
乱序执行；
较大的单线程缓存；
低延迟缓存层次；
少量但强大的核心；
适合复杂控制流；
适合指针密集访问。

例如树遍历：

while (node != nullptr) {
    if (condition(node)) {
        node = node->left;
    } else {
        node = node->right;
    }
}

不同任务路径复杂且数据量不大时，CPU 可能明显优于 GPU。

四十五、为什么 GPU 仍然可以做图计算

GPU 不擅长并不代表不能做。

当图足够大时：

边数量巨大；
顶点数量巨大；
仍然存在大量并行性。

通过专门优化：

CSR/COO 混合格式；
行长度分类；
动态工作队列；
Push/Pull 切换；
Warp 协作；
原子聚合；
数据重排序；

GPU 仍可能获得很高总体吞吐率。

只是它通常难以达到稠密 GEMM 那样的理论峰值利用率。

四十六、原子聚合优化

假设一个 Warp 中多个线程都要更新同一个地址。

朴素方式：

每个 Lane 单独 atomicAdd

可以先在 Warp 内合并相同目标：

识别目标地址相同的 Lane
        ↓
Warp 内先求和
        ↓
只让一个 Lane 执行 Atomic

例如8个线程更新同一地址：

8次 Atomic
    ↓
Warp 聚合
    ↓
1次 Atomic

这称为 Warp-Aggregated Atomics。

它可以显著减少热点原子请求。

四十七、数据重排序

如果列索引或图顶点顺序完全随机，Cache 局部性很差。

可以重新排列：

矩阵行；
矩阵列；
图顶点编号；
边顺序；
Token/Expert 顺序。

目标是让相邻线程访问更接近的数据。

例如将高关联顶点放在相邻编号区域，可以提高：

Cache 命中；
访存合并机会；
分区负载均衡。

但重排序有预处理成本，还需要维护新旧编号映射。

四十八、稀疏格式转换

算法的不同阶段可能偏好不同格式：

构建：COO
SpMV：CSR
按列计算：CSC
Tensor Core：Block Sparse

格式转换可能包括：

排序；
Prefix Sum；
重排；
去重；
合并重复坐标；
构建行指针。

如果只执行一次计算，格式转换成本可能超过计算本身。

如果同一个矩阵重复使用很多次，转换成本可以被摊薄。

四十九、从 ASIC 角度看稀疏加速

ASIC 可以针对特定稀疏格式设计：

Compressed Value Buffer
Index Decoder
Nonzero Scheduler
Gather Unit
MAC Array
Accumulator Router
Scatter / Atomic Unit

核心挑战包括：

如何持续向 MAC 阵列供给非零数据；
如何避免部分 PE 空闲；
如何处理不同长度行；
如何路由激活数据；
如何合并相同输出地址；
如何控制索引和元数据开销。

五十、稀疏 PE 阵列的利用率问题

假设有64个 MAC 单元，但当前周期只有20个有效非零乘法：

\text{Utilization} = \frac{20}{64} =31.25\%

虽然跳过了零乘法，但大量硬件仍然空闲。

如果要提高利用率，需要把来自不同：

行；
通道；
Token；
输出位置；

的非零任务动态组合到同一阵列中。

这会增加调度和路由复杂度。

五十一、稀疏加速器中的元数据

每个非零元素可能需要：

列索引
行边界
Block 编号
偏移
有效掩码

元数据会消耗：

存储容量；
内存带宽；
解码逻辑；
功耗；
时钟周期。

所以评价稀疏加速器时，不能只计算非零 MAC 数量。

还需要统计：

\text{Value Bytes} + \text{Metadata Bytes}

以及索引处理开销。

五十二、结构化稀疏对 ASIC 的价值

固定 N:M 或 Block Sparse 可以使硬件保持规则节拍。

例如2:4稀疏：

每4个权重
读取2个非零值
读取一个小型位置编码
执行2个乘法

控制逻辑固定，容易流水化。

相比完全非结构化稀疏：

每个周期非零数量不确定
位置不确定
输出地址不确定

结构化稀疏通常更容易获得稳定利用率和时序。

五十三、选择稀疏还是低比特量化

稀疏和量化都可以减少模型存储。

例如FP16权重：

16 bit

50%非结构化稀疏后，理论数值数据降为：

8 bit/原始元素

但还要加索引。

INT8稠密量化：

8 bit/元素

不需要逐元素稀疏索引，数据流仍然规则。

因此在某些硬件上：

稠密 INT8 可能比稀疏 FP16 更简单、更快。

更低的INT4同样可能比非结构化稀疏更有吸引力。

五十四、稀疏与量化可以结合

例如：

2:4 结构化稀疏
+
INT8 或 FP8

可以同时减少：

非零数量；
单个数值位宽。

但硬件需要同时支持：

稀疏元数据；
低精度解码；
Scale；
稀疏 Tensor Core 或专用阵列。

组合优化的理论收益很大，但软件和硬件复杂度也更高。

五十五、分析稀疏 Kernel 时看什么

使用 Nsight Compute 时，重点关注：

DRAM / L2 Throughput
Global Load Efficiency
Long Scoreboard
Eligible Warps
Branch Efficiency
Warp Execution Efficiency
Atomic Throughput
Load/Store 指令比例
Occupancy
线程块执行时间差异

典型现象可能是：

计算吞吐率低
显存吞吐率也没有跑满
Long Scoreboard 高
Eligible Warps 少

这通常说明不是简单的峰值带宽受限，而是随机访存造成的延迟受限。

五十六、如何区分带宽受限和延迟受限

带宽受限

表现为：

显存吞吐率接近上限
大量请求连续
Eligible Warp 仍然充足

优化重点：

减少字节数；
压缩数据；
提高复用；
降低精度。

延迟受限

表现为：

显存吞吐率不高
Long Scoreboard 很高
Eligible Warp 很少
访问随机、事务细碎

优化重点：

改善局部性；
重排序；
合并任务；
增加 Memory-Level Parallelism；
减少指针依赖；
使用缓存友好格式。

五十七、本课核心结论

第一，稀疏计算通过跳过零元素减少算术量，但会引入索引、控制和不规则访存开销。

第二，COO 保存每个非零元素的行、列和值，结构简单但索引开销较高。

第三，CSR 按行压缩，适合逐行计算和 SpMV；CSC 按列压缩。

第四，SpMV 每个非零元素只做一次乘加，算术强度很低，通常受到内存系统限制。

第五，x[col_idx] 是典型 Gather 访问，可能导致访存不合并和 Cache 局部性差。

第六，Scatter 更新可能产生写冲突，需要 Atomic。

第七，稀疏矩阵每行非零元素数量不同，会造成线程和 Warp 的负载不均衡。

第八，短行、中等行和长行通常需要不同的线程映射策略。

第九，SpMM 每个非零元素可以参与多个输出通道计算，因此比 SpMV 更容易提高算术强度。

第十，非结构化稀疏自由度高，但硬件和访存效率较差。

第十一，Block Sparse 和 N:M 稀疏通过限制稀疏模式，提高数据流规则性和硬件利用率。

第十二，50%稀疏不等于端到端2倍加速，因为仍有元数据、访存和其他算子开销。

第十三，激活稀疏比固定权重稀疏更难处理，因为稀疏位置在运行时变化。

第十四，图计算的主要困难包括随机邻接访问、顶点度数不均、原子冲突和动态 Frontier。

第十五，GPU 最擅长规则、大规模、控制流一致、访存连续的数据并行任务。

第十六，GPU 不擅长强串行依赖、随机指针访问、复杂分支、热点原子和小规模低延迟任务。

第十七，稀疏计算是否有效，不能只看稀疏率，还要看：

\boxed{ \text{nonzero distribution} + \text{index overhead} + \text{access regularity} + \text{load balance} + \text{hardware support} }

下一课是倒数第二课：

第二十六课：如何从零设计一个简化的 SIMT GPU

将从硬件设计角度把前面知识组合起来：

最小指令集
线程上下文
Warp 与 Active Mask
Warp PC
Block Dispatcher
Register File
Scoreboard
Warp Scheduler
SIMD 执行单元
Load/Store Unit
Shared Memory
Barrier
一个 Kernel 从启动到结束的完整流程
哪些部分可以先用 RTL 实现