CUDA 编程系列之三：内存墙与分块（tiling）的诞生

一张 H100 显卡的纸面算力是 66.9 TFLOPS（FP32）。

现在请它跑一个最简单的程序——向量加法 C[i] = A[i] + B[i]——然后测量它实际达到了多少 TFLOPS。

你会得到一个让人怀疑人生的数字：个位数都不到。用直觉说，就是这张卡只发挥了百分之几的力气。

不是代码写错了，也不是驱动有问题。挡住它的，是 GPU 计算里最著名的一堵墙。PMPP 第五章整章都在拆解这堵墙，并给出了全书最重要的那个优化技术。

1. 两个天花板

任何处理器都有两个性能上限：

上限	含义	H100 的数值
峰值计算吞吐	单位时间内最多做多少次运算	66.9 TFLOPS
峰值内存带宽	单位时间内最多搬多少字节	3.35 TB/s

注意它们量纲不同：一个是"每秒多少次运算"，一个是"每秒多少字节"。

而一个 kernel 的运行速度，恰恰取决于它每搬一个字节，能顺带做多少次运算。这个比率太低，那么无论算力多强，处理器都在等数据——这就是"内存墙"。

2. 算术强度：一个比值定乾坤

这个比值有个正式名字：算术强度（arithmetic intensity），也叫计算访存比，单位是 FLOP/Byte。

$$\text{Arithmetic Intensity} = \frac{\text{FLOPs performed}}{\text{bytes accessed from global memory}}$$

判断规则非常干脆：

比值高 → kernel 是 计算受限（compute-bound），速度由算力决定；
比值低 → kernel 是 带宽受限（memory-bound），速度由带宽决定。

那么分界线在哪里？把两个天花板相除就知道了：

$$\text{threshold} = \frac{66.9 \times 10^{12}}{3.35 \times 10^{12}} \approx 20\ \text{FLOP/B}$$

在 H100 上：算术强度高于 20 FLOP/B 的 kernel 大概率受算力限制，低于 20 的大概率受带宽限制。（换一张卡，这个阈值就变，因为它的算力和带宽之比变了。）

3. Roofline 模型：一张图看清你卡在哪

把这个判断画成图，就是 Roofline 模型：

横轴：算术强度（FLOP/B）；
纵轴：实际达到的计算吞吐（GFLOPS）；
图上有两条线构成"屋顶"：一条由峰值带宽决定的斜线，一条由峰值算力决定的水平线。

两条线的交点，就是那个阈值。斜线以左是"带宽受限区"，水平线以右是"算力受限区"。

你把自己的 kernel 作为一个点画上去，会得到三种截然不同的结论：

点的位置	含义	该做什么
落在斜线上	已经吃满带宽，受强度限制	提高算术强度（向右移动）
落在斜线下方	带宽没用满	改善访存效率（向上移动）
落在水平线上	已经吃满算力	停下来，没有优化空间了

这个模型最大的价值不是算得多准，而是逼你先回答问题再动手：

graph LR A["测量 kernel"] --> B{"算术强度
高于阈值?"} B -- 高于 --> C["计算受限
优化指令吞吐"] B -- 低于 --> D["带宽受限
提高强度 / 改善访存"] C --> E["接近屋顶就停手"] D --> E

4. 算一笔账：向量加法的 0.083

现在回到开头那个让人沮丧的测量结果，把账算清楚。

向量加法里，每个线程做一次浮点加法，同时访问全局内存三次数：

读 A[i]：4 字节
读 B[i]：4 字节
写 C[i]：4 字节

合计 12 字节，完成 1 次浮点运算，所以：

$$\text{Arithmetic Intensity} = \frac{1\ \text{FLOP}}{12\ \text{B}} \approx 0.083\ \text{FLOP/B}$$

对比阈值 20 FLOP/B——差了 240 倍。

这个数字一次性解释了两件事：

为什么向量加法永远跑不到峰值算力。不是代码问题，是这个问题本身"每搬 3 个字的代价只换 1 次加法"。
为什么提高强度的唯一办法是少搬数据。对 C = A + B 来说，三个字是刚需，搬不动了。所以对这类问题，工程上的正确答案往往是"别把它当计算问题看"——它的价值在于它是内存的标尺，而不是算力的标尺。

而矩阵乘法不一样。它有大量可复用的数据，这正是分块技术的用武之地。

5. CUDA 的存储层级

在讲分块之前，先把 GPU 的"家具"摆出来：

存储	位置	作用域	速度	谁管理
寄存器	片上	每线程	最快	编译器 / 程序员
共享内存	片上	每线程块	很快	程序员
常量内存	片上缓存	只读，广播友好	同址访问时很快	程序员
全局内存	片外 DRAM	整个网格	慢（但总带宽高）	程序员

关键点是中间那两行：共享内存和寄存器是程序员可以主动使用的"临时工作台"。它们很小（这是它们快的原因），但正因为小，才需要你精心安排——把要反复使用的数据搬上来，用完再换下一批。

6. 分块：一次搬运，反复取用

回到第一篇留下的那个悬念。朴素的矩阵乘法 P[row][col] = Σ M[row][k] × N[k][col]，每个输出元素要读 $2N$ 个输入，整个矩阵开出了 $2N^3$ 次全局读取。

但请想一想：这些读取里，有多少是重复的？

答案是绝大部分。M 的同一行，会被这一行对应的所有输出元素反复使用；N 的同一列也一样。数据明明就那么点，却被搬了 $N$ 遍。

**分块（tiling）**的思路就在这里：

与其每个线程各搬各的，不如让整个线程块协作，把一块输入搬运到共享内存里；等它上来了，所有线程反复从共享内存取用；用完了，再一起搬下一块。

这就是"分块"这个名字的由来——把大矩阵切成一格一格的小块，块内数据一次搬运、多次使用。

#define TILE_WIDTH 16
__global__ void matrixMulKernel(float* M, float* N, float* P, int Width) {
    __shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
    __shared__ float Nds[TILE_WIDTH][TILE_WIDTH];
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int Row = by * TILE_WIDTH + ty;
    int Col = bx * TILE_WIDTH + tx;
    float Pvalue = 0.0f;                       // 寄存器里累加
    for (int m = 0; m < Width / TILE_WIDTH; ++m) {
        Mds[ty][tx] = M[Row * Width + (m * TILE_WIDTH + tx)];
        Nds[ty][tx] = N[(m * TILE_WIDTH + ty) * Width + Col];
        __syncthreads();                       // 块内数据搬运完毕
        for (int k = 0; k < TILE_WIDTH; ++k)
            Pvalue += Mds[ty][k] * Nds[k][tx];
        __syncthreads();                       // 保护 Mds/Nds，防止被下一轮覆盖
    }
    P[Row * Width + Col] = Pvalue;
}

对比一下这段代码和第一篇的基线版本，三个变化很关键：

共享内存出现了：Mds、Nds 是块内协作的工作台。
累加器挪到了寄存器：Pvalue 在整个块循环期间待在寄存器里，最后才写回全局内存——这叫寄存器分块。
出现了两次 __syncthreads()：一次确认"数据都搬上来了"，一次保护"这轮数据别被下轮覆盖"。

那么账目改善了多少？

每一次全局读取，服务的不再是 1 次使用，而是 $TILE\_WIDTH$ 次使用。全局读取总量从 $2N^3$ 降到大约：

$$ \frac{2N^3}{TILE\_WIDTH} $$

取 $TILE\_WIDTH = 16$，访存量直接降到原来的十六分之一。算术强度被抬高了 16 倍 —— 这正是第四节里那个"向右移动"的动作。

7. 分块的代价：天下没有白吃的午餐

分块看起来完美，但它有三个必须承认的成本。

第一，__syncthreads() 是内在成本，不是附带成本。

分块的"物理意义"是：让一群线程在同一段时间内聚焦于同一小块数据。而"让大家在同一时刻对齐"这件事，本身就需要屏障。所以屏障不是实现细节，它是分块定义的一部分——只要做分块，就一定会引入同步开销。

第二，当输入块和输出块形状不一致时，会出现 halo（边界重叠）。

这在卷积、模板计算（stencil）里尤其明显：算一个 $16\times16$ 的输出块，可能需要 $18\times18$ 的输入块——多出来的那一圈边界数据，被相邻的块重复搬运了。它让"分块到底值不值"变成了一个要算的账：如果块太小、滤波核太大，halo 占比就会高到吞掉全部收益。（这是第七、八章的主题。）

第三，共享内存是从占用率里"抠"出来的。

还记得上一篇的算例吗？把 SM 卡在 75% 的不是线程，也不是寄存器，而是共享内存。你每多申请一个字节的共享内存，就少一份给其他块用的空间。

于是有了这条相当反直觉的规则：

如果一个 kernel 的瓶颈本来就是占用率、而且占用率正是被共享内存卡住的，那么再加分块，会让它更慢。

优化最忌讳的就是"听说这个技术好就用"。你得先知道自己的瓶颈在哪——这也是第六、九、十二章反复要做的事。

顺便补一个小坑：分块版本的边界检查比基线版本更严格。当矩阵边长不是 TILE_WIDTH 的整数倍时，越界的元素必须填 0（而不是跳过），否则它会污染乘加累加的结果。

8. 分块不是 GPU 的专利

最后强调一句 PMPP 特意点明的话：分块是一种通用的并行计算策略，不是 CUDA 的小技巧。

原因很简单：任何并行系统要利用高速缓存，都要求程序具备数据局部性。多核 CPU 上，分块让程序有效利用片上 cache；GPU 上，分块让程序有效利用共享内存。硬件换了一茬又一茬，物理约束没变：小而快的存储总是有限的，只有重复使用才能弥补它的容量。

所以这一章学到的不是"CUDA 有一个叫 shared memory 的东西"，而是：

先问算术强度，再谈优化；而要抬高算术强度，最有力的手段就是——把重复的数据搬到离计算更近的地方。

9. 小结

处理器有两个天花板：峰值算力与峰值带宽，量纲不同，不能直接比。
**算术强度（FLOP/Byte）**决定 kernel 受哪一个限制；阈值 = 峰值算力 ÷ 峰值带宽（H100 约 20 FLOP/B）。
Roofline 模型把"慢"翻译成三种可执行的动作：向右（提高强度）、向上（改善访存）、停下（已到屋顶）。
向量加法 1 FLOP / 12 B = 0.083 FLOP/B，深陷带宽受限——它衡量的是内存，不是算力。
CUDA 提供了寄存器、共享内存、常量内存这些"片上工作台"，但容量有限，需要程序员精心调度。
分块把重复的全局读取变成"一次搬运 + 反复取用"，让朴素矩阵乘法的 $2N^3$ 次读取降到约 $2N^3/TILE\_WIDTH$。
分块的代价是屏障同步、halo 重叠开销，以及被吃掉的占用率——用之前先确认瓶颈。

到这里，“线程如何映射数据”、“硬件如何调度线程”、“数据如何避免反复搬运"这三块拼图已经凑齐了。但还有一个悬而未决的问题：为什么有时候明明搬得不多，带宽还是用不满？

下一篇，我们进入访存本身——合并访问（coalescing）与共享内存的 32 个 bank。

参考链接

Programming Massively Parallel Processors: A Hands-on Approach (5th Edition) — Wen-mei W. Hwu, David B. Kirk, Izzat El Hajj. 第 5 章 Memory architecture and data locality.

1. 两个天花板#

2. 算术强度：一个比值定乾坤#

3. Roofline 模型：一张图看清你卡在哪#

4. 算一笔账：向量加法的 0.083#

5. CUDA 的存储层级#

6. 分块：一次搬运，反复取用#

7. 分块的代价：天下没有白吃的午餐#

8. 分块不是 GPU 的专利#

9. 小结#