第 4 讲 · GPU 架构与 CUDA 编程（二）：矩阵乘法与并行归约 · GPU Architecture & CUDA Programming (2)

对应 Maps to：CMU 15-779（Fall 2025，Zhihao Jia），Lecture 4 · Wed 9/10 · GPU Architecture and CUDA Programming (2)。 阅读材料：原课这一讲没有指定必读论文。推荐：M. Harris, Optimizing Parallel Reduction in CUDA（NVIDIA，2007）；S. Boehm, How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance（博客，2022）；NVIDIA CUDA C++ Best Practices Guide 的访存优化一章。 依据说明：原课第 3、4 讲共用一份 slides（03-CUDA-programming.pdf）。本讲对应其后半部分（第 48–69 页）：案例 1 矩阵乘法（朴素实现、线程级寄存器分块、线程块级共享内存分块、访存复用分析、合作加载）；案例 2 并行归约（树形归约、没有全局同步与内核分解、交错寻址与发散的 warp、步长索引、合并访存、顺序寻址）；以及小结（一致的 warp、合并访存、共享内存 bank 冲突、warp 级优化、Tensor Core）。slides 的分块矩阵乘法代码按 $C = A^\top B$（A 以 $K \times N$ 形状存放）的约定书写，本讲沿用这一约定。全部模拟实验、外推估算和 §6 的转置例子是本讲义自己设计的，不代表课堂上讲了什么。 配套代码 Code：code/lecture04/（4 个脚本；共用第 3 讲的 gpusim.py；只需 numpy；每个脚本 ≤ 15 秒） 预计阅读时间：3–4 小时（含运行代码）

0. 本讲导航 · How to read this lecture

0.1 学习目标

学完这一讲，你应该能够：

1. 分析朴素矩阵乘法的访存量（$2N^3$ 次全局读取），并解释为什么线程编号和数据布局的对应关系决定了访存能否合并。
2. 推导线程级寄存器分块（全局访问 $2N^3/V$）和线程块级共享内存分块（全局访问 $2N^3/L$、共享内存访问 $2N^3/V$），理解”两级分块各管一层存储”；会写合作加载。
3. 用三层 roofline（HBM、共享内存、计算）估算分块矩阵乘法的速度，并说出分块参数受哪些资源约束（寄存器、共享内存、占用率）。
4. 解释为什么 CUDA 没有块间全局同步，怎样用内核分解实现全局归约。
5. 比较并行归约的几个版本：交错寻址（warp 发散）、步长索引（bank 冲突）、顺序寻址（两者都没有）、warp shuffle（减少同步）。
6. 区分合并访存（全局内存：一个 warp 的地址要连续）和bank 冲突（共享内存：一个 warp 的地址要分散在不同 bank），会用 padding 消除转置里的 bank 冲突。

0.2 先修与本讲范围

0.3 本讲地图

§1  为什么要拿矩阵乘法和归约做案例
§2  矩阵乘法（一）：朴素实现与"x 当行还是当列"                     ← 代码 01
§3  矩阵乘法（二）：寄存器分块与共享内存分块，复用分析                ← 代码 02（演示 1）
§4  矩阵乘法（三）：合作加载；三层 roofline；分块参数的权衡            ← 代码 02（演示 2–3）
§5  并行归约：内核分解、交错寻址、步长索引、顺序寻址、warp shuffle      ← 代码 03
§6  合并访存与 bank 冲突：跨步、对齐、矩阵转置                        ← 代码 04
§7  回顾：GPU 优化技巧清单
§8  小结   §9 常见误区   §10 练习   §11 延伸阅读   §12 下一讲预告

0.4 一个提醒：本讲的计数是怎么来的

和第 3 讲一样，所有计数都来自 gpusim.py：每个内核在 $N = 256$（矩阵乘法）或 $N = 2^{18}$（归约）上真实执行一遍，结果和 NumPy 对照，同时数出每条 warp 级访存指令的 32 字节事务数、共享内存的 wavefront（bank 冲突时一条指令要走几轮）、发散次数等。然后按”读流量 ∝ $N^3$、写流量 ∝ $N^2$“（矩阵乘法）或”∝ $N$“（归约）外推到真实规模，用 H100 的标称值估算时间：HBM 3.35 TB/s；共享内存每个 SM 每周期 1 个 wavefront（128 字节）；CUDA core FP32 66.9 TFLOP/s（本讲的内核都不用 Tensor Core）。这些时间是”按瓶颈估算的下界”，不是实测；它们的价值在于比较版本之间的差别和判断瓶颈。

本讲的模拟器新增了向量访存：t.load(A, idx, vec=True) 表示一个线程用一条指令读 $V$ 个连续的元素（对应 CUDA 的 float4）。真实的高性能内核都这样读写，它减少了指令数。

1. 为什么拿矩阵乘法和归约做案例 · Why these two case studies

slides 选的两个案例代表了深度学习里最重要的两类计算：

• 矩阵乘法：Transformer 里 99% 以上的 FLOPs（第 1 讲 §8：预填充时矩阵乘法占 99.97% 的 FLOPs）。它是计算密集的——但只有在数据被充分复用时才是。怎样组织复用，是本讲的主线。
• 归约（求和、求最大值）：softmax、LayerNorm、RMSNorm、注意力里的归一化都离不开它。它是访存密集的——每个数只用一次，优化的目标是”不要比读一遍数据更慢”。

两个案例覆盖了第 3 讲的几乎所有概念：合并访存、共享内存、同步、发散、bank 冲突、占用率。它们也是后面几讲的零件：FlashAttention（第 5 讲）= 分块矩阵乘法 + 在线 softmax（一种归约）；Triton（第 8 讲）把这里手写的分块交给编译器。

2. 矩阵乘法（一）：朴素实现 · Naive matmul

2.1 slides 的”稻草人”

slides 第 49 页：计算 $C = A \times B$（$N \times N$），每个线程算 $C$ 的一个元素，块大小 $32 \times 32$：

__global__ void mm(float A[N][N], float B[N][N], float C[N][N]) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    float result = 0;
    for (int k = 0; k < N; ++k) result += A[x][k] * B[k][y];
    C[x][y] = result;
}

slides 的分析：每个线程读全局内存 $2N$ 次，共 $N^2$ 个线程，总共 $2N^3$ 次。每做一次乘加，就要从显存读两个数——算术强度只有 $2 \text{ FLOP} / 8 \text{ 字节} = 0.25$，比 H100 的拐点低三个数量级。

2.2 “x 当行还是当列”：合并访存的第一课（代码 01）

slides 的这段代码有一个容易被忽略、但影响巨大的细节：x 来自 threadIdx.x，而 threadIdx.x 是一个 warp 里变化最快的那一维（第 3 讲 §8）。于是一个 warp 的 32 个线程有 32 个连续的 x、相同的 y：

• 读 A[x][k]：32 个线程读 A 的同一列的 32 个元素，相邻地址隔 $N$ 个 float——32 个不同的 32 字节段，32 个事务；
• 读 B[k][y]：32 个线程读同一个数——广播，1 个事务；
• 写 C[x][y]：又是一列，32 个事务。

把行列交换（row 来自 threadIdx.y、col 来自 threadIdx.x）：读 A[row][k] 变成广播（1 个事务），读 B[k][col] 是一行里连续的 32 个数（4 个事务），写 C[row][col] 也是连续的（4 个事务）。代码 01 在 $N = 256$ 上的计数：

两个版本访问次数完全相同（都等于 slides 的 $2N^3$），事务数差了 6.6 倍。外推到 $N = 4096$（FP32，H100）：slides 原版约 677 ms，交换版约 103 ms——而纯计算只需要 2.05 ms。

规则：让 threadIdx.x（warp 里变化最快的维度）对应内存里连续的维度（行优先存储时就是列下标）。slides 的代码是为了说明访问次数，真正写内核时这一点比什么都重要。

2.3 复用在哪里？

即使合并得最好，朴素版每次乘加也要从显存搬约 2.5 字节（算术强度 0.40）。可是矩阵乘法里有大量复用：A 的每个元素被 $N$ 个线程（C 的一整行）用到，B 的每个元素也被 $N$ 个线程用到。如果有一个”理想的缓存”让每个数只从显存读一次，算术强度可以到 $2N^3 / (3N^2 \times 4) = N/6 \approx 683$。真 GPU 的 L1 / L2 能挡住一部分重复读，但它们很小（H100 的 L2 50 MB，装不下三个 64 MB 的矩阵），而且替换策略不受程序员控制。分块就是把复用显式地组织起来。

3. 矩阵乘法（二）：两级分块 · Two-level tiling

3.1 优化 1：线程级寄存器分块

slides 第 50 页：让每个线程计算 $C$ 的一个 $V \times V$ 小块。每一步 $k$，线程从全局内存读 A 的一小段（$V$ 个数）和 B 的一小段（$V$ 个数）到寄存器，做一个外积，更新 $V^2$ 个累加器：

float c[V][V] = {0};
float a[V], b[V];
for (int k = 0; k < N; ++k) {
    a[:] = A[k, xbase*V : xbase*V + V];      // 读 V 个数
    b[:] = B[k, ybase*V : ybase*V + V];      // 读 V 个数
    for (y...) for (x...) c[x][y] += a[x] * b[y];   // V² 次乘加
}

为什么有效？ 读进寄存器的每个数被用了 $V$ 次（$a[x]$ 和 $V$ 个 $b[y]$ 相乘）。每个线程读 $2NV$ 次，线程数变成 $N^2/V^2$，全局访问总量 $2N^3/V$。

3.2 优化 2：线程块级共享内存分块

slides 第 51–52 页：再往上一层。让一个线程块计算 $C$ 的一个 $L \times L$ 块，块里每个线程仍然算 $V \times V$。沿 $k$ 维每次推进 $S$：块里的线程合作把 A 和 B 的一个 $S \times L$ 条带搬进共享内存，同步，然后每个线程从共享内存读数据、在寄存器里做外积：

__shared__ float sA[S][L], sB[S][L];
float c[V][V] = {0};
for (int ko = 0; ko < N; ko += S) {
    __syncthreads();
    sA[:, :] = A[ko : ko + S, yblock*L : yblock*L + L];   // 合作加载（§4.1）
    sB[:, :] = B[ko : ko + S, xblock*L : xblock*L + L];
    __syncthreads();
    for (int ki = 0; ki < S; ++ki) {
        a[:] = sA[ki, threadIdx.y*V : threadIdx.y*V + V];
        b[:] = sB[ki, threadIdx.x*V : threadIdx.x*V + V];
        c += outer(a, b);
    }
}

（slides 的这段代码里第二个 sA 应为 sB，是一个小笔误。）注意 A[ko : ko+S, yblock·L : ...] 的写法：这里的 A 以 $K \times N$ 的形状存放，计算的是 $C = A^\top B$——这样 A 和 B 的条带在内存里都是”一行里连续的 $L$ 个数”，合作加载天然合并。代码 02 沿用了这个约定。

复用分析（slides 第 52 页）：

• 全局内存：每个块读 $2LN$ 个数（A、B 各 $L \times N$），块数 $N^2/L^2$，总量 $2N^3/L$。每个从显存读进共享内存的数，被块里的线程用了 $L$ 次。
• 共享内存：每个线程读 $2NV$ 个数，线程数 $N^2/V^2$，总量 $2N^3/V$。每个从共享内存读进寄存器的数，被用了 $V$ 次。

3.3 数一数（代码 02 演示 1）

共享内存分块里，共享内存的读取量是 8,388,608 个元素，正好等于 $2N^3/V$。实际计数和 slides 的公式完全吻合。

两级分块各管一层存储，这是本讲最重要的直觉：

寄存器  ←──(每个数用 V 次)──  共享内存  ←──(每个数用 L 次)──  HBM 显存
         寄存器分块管这一层                共享内存分块管这一层

存储层次的每一层都有自己的带宽上限，每一层都需要自己的复用。只做共享内存分块不做寄存器分块，共享内存的带宽会成为瓶颈；只做寄存器分块，HBM 仍然是瓶颈。

4. 矩阵乘法（三）：合作加载与参数权衡 · Cooperative fetching & trade-offs

4.1 合作加载

slides 第 53 页：搬一个 $S \times L$ 的条带进共享内存，要让块里的所有线程一起搬，并且相邻的线程搬相邻的数：

int nthreads = blockDim.y * blockDim.x;
int tid = threadIdx.y * blockDim.x + threadIdx.x;
for (int j = 0; j < L * S / nthreads; ++j) {
    int y = (j * nthreads + tid) / L;
    int x = (j * nthreads + tid) % L;
    s[y, x] = A[k + y, yblock * L + x];
}

线性编号 tid 连续的线程，得到连续的 x（同一行里相邻的列），读的是连续的地址——合并。代码 02 演示 2 把编排反过来（y = lin % S, x = lin / S，相邻线程搬同一列里相邻的行）作对比：

同样的数据、同样的线程，换一个编排，全局事务翻倍，共享内存的写入也开始冲突。

4.2 三层 roofline：哪组参数最快？（代码 02 演示 3）

把每个配置都在 $N = 256$ 上模拟，再外推到 $N = 4096$（FP32，H100 CUDA core）。每个配置的时间取三者中最大的：计算时间、HBM 时间（事务 × 32 字节 ÷ 3.35 TB/s）、共享内存时间（wavefront ÷ 每 SM 每周期 1 个）：

（寄存器数按”$V^2$ 个累加器 + $2V$ 个 a、b + 约 16 个杂项”粗估。）

读这张表的方法：

1. $L$ 压 HBM，$V$ 压共享内存。 $L$ 从 64 增大到 128，HBM 流量减半；$V$ 从 4 增大到 8，共享内存时间从 1.80 降到 0.64 ms。
2. 压到计算成为瓶颈就够了。 $L=128, S=8, V=8$ 时 HBM 和共享内存都快于计算，估算 2.05 ms = 纯计算时间。这就是第 1 讲说的”把访存受限的运算推到 roofline 的平顶上”。
3. 资源有上限。 $V=16$ 要约 304 个寄存器，超过每线程 255 个的硬上限，编译器会把多出来的”溢出”（spill）到显存里，性能断崖式下降——这个简单模型没有体现。块越大、每块要的共享内存越多，能驻留的块越少。
4. 最快的配置占用率只有 25%。 roofline 不考虑延迟隐藏（第 3 讲 §7.5）。高性能 GEMM 的做法正是”低占用率 + 每个线程大量独立的乘加”（指令级并行）来藏延迟；第 6 讲的异步拷贝和 warp 特化会把”搬数据”和”算”彻底重叠。

没有一组参数对所有矩阵形状、所有 GPU 都最好：换一块共享内存更大的卡、换一个”瘦”矩阵（第 1 讲 §4.4 的 $K=1$ 例子），最优参数都会变。cuBLAS 内部为不同形状准备了大量手调的内核，而第 9 讲的自动调优让机器去搜。

5. 并行归约 · Parallel reduction

5.1 问题与挑战

slides 第 55–57 页：求 $\sum_{i=1}^{n} A[i]$。它是归一化、softmax 等算子的常用原语。块内用树形归约：

3   1   7   0   4   1   6   3
  4       7       5       9
      11              14
              25

挑战在于：为了用满 GPU，必须用很多个线程块（一个块只在一个 SM 上运行），每个块归约数组的一段；那么块与块之间的部分和怎样合并？

CUDA 没有全局同步。 回忆第 3 讲的核心假设：线程块可以按任意顺序执行，块之间不能同步。slides 给了两个理由：

• 为处理器数量很多的 GPU 构建全局同步硬件很昂贵；
• 如果块数超过”SM 数 × 每 SM 可驻留的块数”，先运行的块在栅栏处等待，而还没被调度的块永远等不到空出来的 SM——死锁。

解决办法：内核分解。 把计算拆成多次内核启动：第一个内核让每个块产生一个部分和，第二个内核再归约这些部分和……每一层用的是同一份代码。内核启动充当全局同步（第 3 讲 §6：内核边界是隐式的全局栅栏），而启动开销很低（几微秒）。代码 03 演示 2：$N = 2^{18}$、每块 256 个线程时，v1–v3 要启动 3 个内核（1024 → 4 → 1 个块），每个块处理 512 个数的 v4 只要 2 个（512 → 1）。所有版本的结果和 float64 参考值相差 $10^{-5}$ 量级（float32 累加顺序不同造成的舍入差）。

5.2 版本 1：交错寻址

slides 第 60–62 页：

sdata[tid] = g_idata[i];
__syncthreads();
for (unsigned int s = 1; s < blockDim.x; s *= 2) {
    if (tid % (2 * s) == 0) sdata[tid] += sdata[tid + s];
    __syncthreads();
}
if (tid == 0) g_odata[blockIdx.x] = sdata[0];

问题：warp 高度发散。 slides 列出了每一轮哪些线程是活跃的：

tid:  0 1 2 3 4 5 6 7 8 9 ...
s=1   T F T F T F T F T F ...
s=2   T F F F T F F F T F ...
s=4   T F F F F F F F T F ...

活跃的线程越来越稀疏，但它们分散在所有 warp 里——每个 warp 都要执行，却只有少数车道在干活。

5.3 版本 2：步长索引，不发散的分支

slides 第 63–64 页：把活跃的线程挤到块的前面：

for (unsigned int s = 1; s < blockDim.x; s *= 2) {
    int index = 2 * s * threadIdx.x;
    if (index < blockDim.x) sdata[index] += sdata[index + s];
    __syncthreads();
}

现在活跃的总是编号最小的那些线程（s=1 时前一半，s=2 时前四分之一……），整个 warp 要么全部活跃、要么全部空闲，直到活跃线程少于 32 个。

新问题：bank 冲突。 线程 tid 访问 sdata[2·s·tid]：同一个 warp 里相邻线程的地址相隔 $2s$ 个字。共享内存有 32 个 bank，每个 bank 4 字节宽，连续的字依次落在 bank 0、1、……、31、0……；相隔 $2s$ 个字的 32 个地址只用到 $32 / (2s)$ 个 bank（$s = 1$ 时 16 个、$s = 2$ 时 8 个……），同一个 bank 的不同字要排队——$2s$ 路冲突。

5.4 版本 3：顺序寻址

slides 第 68 页：把循环反过来，步长从大到小：

for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (threadIdx.x < s) sdata[threadIdx.x] += sdata[threadIdx.x + s];
    __syncthreads();
}

活跃的线程仍是编号最小的那些（不发散），而且同一个 warp 访问连续的 32 个字，落在 32 个不同的 bank——没有冲突。slides 把这一页叫做”完全合并的访问”，对共享内存来说更准确的说法是”无 bank 冲突”（合并是全局内存的概念，§6）。

5.5 版本 4：warp 级优化

slides 的小结列出了”warp 级优化”。代码 03 加了一个版本 4，组合两个经典技巧：

1. 加载时先两两相加：每个线程从全局内存读两个数、先加起来再存进共享内存。一个块处理 $2 \times$ 块大小 个数，块数减半，第一轮”一半线程闲着”的浪费也没了。
2. 最后一个 warp 用 shuffle：剩下 64 个数时，只让第一个 warp 干活：每个线程先读 sdata[tid] + sdata[tid+32] 到寄存器，再用 __shfl_down_sync 在寄存器之间对折 5 次（16、8、4、2、1）。同一个 warp 的线程在硬件上一起执行，交换寄存器不需要经过共享内存，也不需要 __syncthreads()。

5.6 数一数（代码 03 演示 1、3）

只看第一个内核（$N = 2^{18}$，1024 个块 × 256 个线程；v4 是 512 个块）：

外推到 $N = 2^{26}$（约 6700 万个数，268 MB），按”指令发射（每 SM 4 个调度器）、共享内存（每 SM 每周期 1 个 wavefront）、HBM”三者中最慢的估算，再加上内核启动开销：

这张表讲了一个完整的故事：

• v1 → v2：消除发散，发射的指令减半（69.7 → 34.6 µs）。但共享内存的总轮数一点没少（v1 是”很多条指令、每条 1 轮”，v2 是”更少的指令、每条 3.33 轮”），瓶颈只是从”发散”换成了”bank 冲突”，估算时间不变。只修一个瓶颈，可能什么也得不到。
• v2 → v3：消除 bank 冲突，共享内存时间降到 45 µs，第一次让 HBM（读一遍数据，80 µs）成为瓶颈。
• v3 → v4：指令和共享内存进一步降低，已经远低于 HBM；剩下的差别来自少一次内核启动。归约每个数只做一次加法，理想情况就是纯访存受限（80 µs）——优化的终点是”和读一遍数据一样快”。

5.7 “为什么需要两个 __syncthreads()？”

slides 第 60 页留了一个问题。代码 03 演示 4 把它们分别删掉：

• 循环里的：保证下一轮读到的是这一轮写完的值。删掉后竞争检测立刻报警。
• 加载之后的：保证所有数都已经搬进共享内存。在 v1 里删掉它，模拟器没有报警——原因是 v1 的第一轮（$s = 1$）每个线程只读同一个 warp 里邻居的数，跨 warp 的读取要到 $s \ge 32$ 才出现，而那时循环里的 __syncthreads() 早已执行过。但这不代表可以删：从 Volta 架构起，同一个 warp 的线程也不再保证同步执行（independent thread scheduling），warp 内通过共享内存交换数据同样需要同步（__syncwarp()）——模拟器只检查跨 warp 的竞争，抓不到这一类。同样的删法放在 v3 里（第一轮就读 sdata[tid + 128]，是别的 warp 写的），竞争 131,072 次。

教训：同一条同步规则是否”恰好”被违反，取决于寻址方式。要按规范写同步，而不是靠测试碰运气。

6. 合并访存与 bank 冲突 · Coalescing & bank conflicts

6.1 合并访存（全局内存）

slides 第 65 页：当一个 warp 的多个线程访问连续的地址时，硬件把它们合并成少数几个事务，带宽利用率最高；地址分散时，同样多的数据需要更多事务。现代 NVIDIA GPU 的全局内存事务粒度是 32 字节（一个 sector）。代码 04 演示 1：一个 warp 的 32 个线程各读一个 float（共 128 字节）：

跨步 ≥ 8 时，每个线程独占一个 32 字节段，只用到其中 4 字节——这正是 slides 原版矩阵乘法读 A 的情形（§2.2）。对齐也有代价：错开一个元素就多搬 25%（第 3 讲的一维卷积里也出现过）。

6.2 bank 冲突（共享内存）

共享内存被分成 32 个 bank，每个 bank 每个周期提供 4 字节。一个 warp 的一条共享内存指令：

• 32 个线程访问 32 个不同的 bank → 1 轮（wavefront）完成；
• 多个线程访问同一个字 → 广播，不算冲突；
• 多个线程访问同一个 bank 的不同字 → 要排队，排得最长的那个 bank 决定要走几轮。

注意两者的”好模式”恰好不同：全局内存要连续，共享内存要分散在不同的 bank（连续的 32 个字恰好满足两者，所以大多数时候不冲突；问题出在”按列访问”）。

6.3 一个综合例子：矩阵转置（代码 04 演示 2）

转置 out[x][y] = in[y][x] 必然有一边按列访问：

1. 朴素版：按行读（4 个事务），按列写（32 个事务）。
2. 共享内存分块：每个块先把 $32 \times 32$ 的子块按行读进共享内存 tile[32][32]，同步，然后按列从共享内存取出（tile[tx][ty]）、按行写出去。读写全局内存都只要 4 个事务了——但按列读共享内存时，同一个 warp 的 32 个线程访问的地址相隔 32 个字，全部落在同一个 bank：32 路冲突，每条读指令要走 32 轮。
3. padding：把共享内存声明成 tile[32][33]，每行多留一个元素。按列读的地址相隔 33 个字，正好错开到 32 个不同的 bank，冲突消失。代价只是 128 字节的共享内存。

又是”只修一个瓶颈不够”：分块解决了全局内存的合并问题，却引入了共享内存的冲突，直到 padding 才把两边都理顺。

7. 回顾：GPU 优化技巧清单 · Recap

slides 最后一页列出的 GPU 优化技巧，正好是第 3–4 讲的内容：

再加上本讲的主线——分块复用（每一层存储都需要自己的复用）和占用率 / 延迟隐藏（第 3 讲 §7）——就是手写高性能 CUDA 内核的基本功。下一个问题自然是：这么多细节，能不能交给编译器？第 8–9 讲的 Triton、TVM、Ansor 就是在回答它。而在那之前，第 5 讲先用这些基本功拆解一个真实的、改变了 LLM 训练的内核：FlashAttention。

8. 小结 · Summary

1. 朴素矩阵乘法读全局内存 $2N^3$ 次。线程编号与数据布局的对应决定能否合并：slides 原版（threadIdx.x 当行号）每条读指令 32 个事务，交换后 1–4 个，事务总数差 6.6 倍。
2. 寄存器分块把全局访问降到 $2N^3/V$，共享内存分块降到 $2N^3/L$、共享内存访问 $2N^3/V$——模拟计数和 slides 的公式完全吻合。两级分块各管一层存储。
3. 合作加载要让相邻线程搬相邻的数；编排反过来，事务翻倍。
4. 三层 roofline：$L$ 压 HBM、$V$ 压共享内存，直到计算成为瓶颈（$N=4096$ 时 $L=128, V=8$ 估算 2.05 ms，比朴素快约 50 倍）；但受寄存器、共享内存和占用率约束，最优参数随硬件和形状变化。
5. 并行归约：没有全局同步 → 内核分解；交错寻址发散（SIMD 效率 43%），步长索引把发散换成了 bank 冲突（3.33 轮/指令、估算时间不变），顺序寻址两者都消除（第一次达到访存受限），warp shuffle 再减少同步和指令。
6. 合并访存 vs bank 冲突：前者要连续，后者要分散；转置用分块解决前者、用 padding 解决后者。只修一个瓶颈往往不够。

9. 常见误区 · Common misconceptions

1. “访问次数一样，速度就一样。” §2.2：同样 $2N^3$ 次访问，事务数差 6.6 倍。决定速度的是事务数（字节），不是访问次数。
2. “共享内存分块就够了。” 不做寄存器分块，共享内存带宽会成为新的瓶颈（§4.2：$V$ 从 4 到 8，共享内存时间降到约三分之一）。
3. “分块越大越好。” 寄存器溢出、共享内存不够、占用率太低、尾部效应（第 3 讲 §7.6）都会反过来伤害性能。
4. “消除了发散就快了。” §5.6 的 v2：发散没了，bank 冲突来了，估算时间不变。优化要盯着当前的瓶颈。
5. “合并访存和无 bank 冲突是一回事。” 一个针对全局内存（要连续），一个针对共享内存（要分散到不同 bank）；按列读共享内存的子块正是”合并了却冲突”的典型。
6. “同一个 warp 的线程天然同步，warp 内交换数据不用同步。” Volta 以后不再成立；warp 内通过共享内存通信要 __syncwarp()，shuffle 指令要带上参与线程的掩码（__shfl_down_sync(mask, ...)）。
7. “块之间可以用一个全局变量自旋等待来同步。” 块数超过可驻留的块数时会死锁；需要全局同步就拆成多个内核（或者用协作组 cooperative groups 的网格同步，它要求块数不超过可驻留的上限）。

10. 练习 · Exercises（含参考答案）

练习 1（朴素矩阵乘法的事务）：在 slides 原版中，如果块大小改成 $(16, 16)$，读 A 和读 B 每条 warp 指令各要几个事务？

答：一个 warp 现在是 16 个 x × 2 个 y。读 A[x][k]：16 个不同的行，每行一个数，16 个事务（两个 y 读同样的地址，广播）。读 B[k][y]：2 个不同的数，可能落在同一个或相邻的段，1 个事务。仍然不合并，只是比 32 好一些。

练习 2（寄存器分块的复用）：$V = 8$ 时，每个线程每一步 $k$ 读多少个数、做多少次乘加？全局访问总量比朴素版少几倍？

答：读 $2V = 16$ 个数，做 $V^2 = 64$ 次乘加；全局访问 $2N^3/8$，少 8 倍。

练习 3（共享内存分块的资源）：$L = 128, S = 16, V = 8$。每块多少线程？每块要多少共享内存？在 H100 SM 上（228 KB 共享内存、64 个 warp、65536 个寄存器）、每线程 96 个寄存器时能驻留几个块？

答：线程 $(128/8)^2 = 256$；共享内存 $2 \times 16 \times 128 \times 4 = 16$ KB；寄存器限制 $65536 / (96 \times 256) = 2.67 \to 2$ 块；共享内存限制 $228/16 = 14$ 块；warp 限制 $64/8 = 8$ 块。驻留 2 块，占用率 $16/64 = 25\%$，受寄存器限制（和代码 02 表中一致）。

练习 4（三层 roofline）：用 FP32 CUDA core 算 $N = 8192$ 的方阵乘法，$L = 128$、$V = 8$。按 §4.2 的模型，HBM 流量大约多少 GB？是计算受限还是 HBM 受限？

答：读 $2N^3/L \times 4$ 字节 $= 2 \times 8192^3 / 128 \times 4 \approx 34.4$ GB，写 C 约 0.27 GB；$t_\text{HBM} \approx 34.6 / 3350 \approx 10.3$ ms；$t_\text{comp} = 2 \times 8192^3 / 66.9\text{T} \approx 16.4$ ms。计算受限（$N$ 越大，同样的分块下算术强度越高）。

练习 5（归约的 bank 冲突）：版本 2 里，$s = 4$ 时一个 warp 的 32 个线程访问 sdata[8·tid]。是几路冲突？

答：相邻线程相隔 8 个字，32 个地址只落在 $32/8 = 4$ 个 bank（0、8、16、24），每个 bank 8 个不同的字：8 路冲突（$2s = 8$）。不过 $s = 4$ 时只有 index < blockDim 的线程活跃——块大小 256 时活跃的是 tid < 32，正好一个 warp。

练习 6（内核分解）：$N = 10^8$，每块 512 个线程、每个线程先加两个数（每块处理 1024 个数）。要启动几个内核？每个内核各多少块？

答：$\lceil 10^8 / 1024 \rceil = 97{,}657$ 块 → $\lceil 97657/1024 \rceil = 96$ 块 → 1 块：3 个内核。

练习 7（padding 为什么是 33）：转置时 tile[32][33] 能消除冲突。如果写成 tile[32][34] 呢？tile[32][64] 呢？

答：按列读的地址相隔 $w$ 个字，落在 bank $(k \cdot w) \bmod 32$。$w = 33$：$k \bmod 32$，32 个不同的 bank，无冲突。$w = 34$：$2k \bmod 32$，只用到 16 个偶数 bank，2 路冲突。$w = 64$：全落在 bank 0，32 路冲突。规律：$w$ 要和 32 互质。

练习 8（读模拟器代码）：在 gpusim.py 里找到计算 bank 冲突的函数，说明它为什么把”同一个 warp 访问同一个字”只算一次。

答：BlockContext._account_shared 先对 (warp, 字地址) 去重（np.unique），这对应硬件的广播：多个线程读同一个字时，bank 只需要提供一次。然后按 (warp, bank) 计数，取每个 warp 里计数最大的 bank 作为这条指令的轮数。

11. 延伸阅读 · Further reading

• M. Harris, Optimizing Parallel Reduction in CUDA（NVIDIA，2007）：本讲 §5 的源头，一共 7 个版本，后几个版本讲了循环展开和模板化。
• S. Boehm, How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance（博客，2022）：从朴素版一步步优化到接近 cuBLAS，和本讲 §2–§4 的思路一致，并给出了 A6000 上的实测数字。
• NVIDIA CUTLASS 文档：Efficient GEMM in CUDA：工业级 GEMM 的层次分块（线程块 / warp / 线程 / Tensor Core 指令）。
• V. Volkov, Better Performance at Lower Occupancy（GTC 2010）：§4.2 那个”占用率 25% 却最快”的现象。
• NVIDIA, CUDA C++ Best Practices Guide，“Memory Optimizations”一章：合并访存、共享内存 bank、对齐的官方说明。
• 本仓库 CMU11-868 第 4 讲：同样的分块矩阵乘法与转置，侧重 miniTorch 作业的写法。

12. 下一讲预告 · Next lecture

第 5 讲《案例：Transformer、注意力与 FlashAttention》把本讲的两个零件拼起来：注意力 = 两次矩阵乘法 + 一次按行的 softmax（归约）。朴素实现要把 $N \times N$ 的注意力矩阵写回显存再读出来，是严重的访存受限；FlashAttention 用分块 + 在线 softmax，让整个注意力在片上完成，并在反向传播时用重计算（第 2 讲 §7）代替存储。我们还会看到，推理时的解码阶段为什么需要另一种切分方式（Flash-Decoding）。第 5 讲有两篇必读论文：Attention Is All You Need 和 FlashAttention。