原文：第 5 章说明：忠实翻译原网页内容，并补入与经典文献、业界系统的对照。术语首次出现给出英文锚点。

第 5 章 CUDA 编程作为软硬件协同优化

从朴素矩阵乘法到分层分块

本章在体系中的位置

第 3 章带你钻进 GPU 内部（SM、warp、HBM、共享内存、寄存器、张量核），第 4 章把 CUDA 编程模型（kernel、block、grid、threadIdx）摆到你面前。但「知道这些概念」和「写出能打的内核」之间隔着一条很宽的鸿沟：为什么一段数学上完全正确的矩阵乘法，在 H100 上只能跑出理论峰值百分之一都不到的吞吐？本章就用 GEMM 这一个具体案例，把第 3、4 章积累的架构约束从头到尾实地走一遍——从 DRAM 的物理组织、warp 的锁步执行，到共享内存与寄存器的分层重用。全章的优化阶梯（合并访问 → 共享内存平铺 → 寄存器平铺 → 向量化 → warp 平铺 → 存储体冲突规避）既是 CUDA 手工优化的巅峰实践，也是第 7 章编译器话题的直接引子：当你看到这一整条阶梯必须在布局、线程映射、平铺尺寸、向量宽度、bank 映射等多个维度上同时调对，你自然会问——这些能不能交给自动搜索？能。这正是编译器与自动调优存在的原因。

在第 3、4 章，我们先后研究了 GPU 微架构与 CUDA 编程模型。你已经见过那台机器的计算资源——CUDA 核、张量核；也见过它的存储与通信层次——从 HBM 到片上缓存、共享内存、寄存器。现在要往前走一步：把那些知识真正用起来。

这一章从头到尾贯穿一个判断：CUDA 编程本质上是软硬件协同优化（hardware-software co-optimization）。软件不是「用 CUDA 写下的算法」那么简单，它是算法在 GPU 执行与存储层次上的一种特定映射（mapping）。内核启动语法、__global__ 关键字、块维度、线程索引，都只是表面现象；更深的问题始终是同一个——一个高维工作负载，应该如何被布局和调度，才能让 GPU 硬件以它能够高效执行的形式看到数据和计算？映射对得上硬件，性能可以非常出色；映射很别扭，一个理论上再漂亮的工作负载也能跑得极差。

让这一切变得生动起来的示例是矩阵乘法。它为什么是机器学习系统课程里最合适的主角？因为它有两个看似矛盾的性质：一方面，它在数学上极其简单，理论算术强度（arithmetic intensity）也漂亮；另一方面，一旦你真在 GPU 上实现它，几乎每一个重要的系统问题都会立刻暴露——数据布局错配、执行错配、合并访问、数据重用、存储放置、存储体冲突、局部性与并行度之间的权衡。矩阵乘法不只是内核，它是一台显微镜，透过它可以看到 CUDA 编程的全部肌理。

1. CUDA 编程首先是映射问题

直觉。 机器学习里的工作负载通常是一个张量计算。矩阵是二维张量，一批矩阵是高维张量；卷积和 transformer 的张量往往带有更丰富的逻辑结构。这些张量的邻接关系天然是多维的：矩阵元素有行邻居和列邻居，图像像素有空间邻居，批量张量里的局部性可能同时横跨 batch、channel、height、width 四个维度。

关键约束。 硬件不会直接「看到」这种逻辑结构。内存是线性编址的——在最底层，DRAM 不是「二维数学」，而是按行、列、页、突发传输组织起来的可寻址存储介质；共享内存是容量很小、由软件显式管理的片上 SRAM；寄存器是线程私有的；GPU 调度器也不会把二维网格当作抽象的几何对象来执行，它只执行一维的 warp——32 条通道锁步（lockstep）前进。

这是 CUDA 里第一重深刻的概念张力：工作负载是高维的，但搬运和执行它的底层机制，几乎是一维的、线性的、分层的。 课程由此指出三类具体的错配：

数据布局错配（data-layout mismatch）。 多维张量必须被放进线性内存。一旦选定某种布局，比如行主序（row-major），你就沿一个维度保住了局部性，同时沿另一个维度削弱了它。
执行错配（execution mismatch）。 你可以用二维或三维线程坐标写代码，但硬件只执行一维 warp。于是从 threadIdx.{x,y,z} 到问题域的映射绝不是语法细节——它决定哪些线程在 warp 里成为邻居，进而决定哪些地址会被一起访问。
通信放置错配（communication placement）。 同样的算术，数据移动发生在哪一层，成本天差地别。重用寄存器里的值极其便宜，重用共享内存里的值也很好（但没那么便宜），反复从 HBM 取数则贵得多。算术没有变，通信基底变了，观察到的性能就完全变了。

记住这句话

CUDA 优化，本质上就是渐进地消除数据布局错配、执行错配和不良的通信放置。抓住这一条，后面一长串技巧就不再是零散口诀，而会呈现出一致的系统逻辑。

这种抽象表述不是哲学空谈，它预测了后面出现的每一个优化。全章 18 个小节，本质上是同一句话的层层展开。

2. 运行示例：一个最朴素的矩阵乘法内核

为了让讨论落地，课程选用方阵乘法：

C = A \times B + C,

其中 $A$ 、 $B$ 、 $C$ 都是按行主序存储的 $N \times N$ 矩阵。这是标准的类 GEMM 操作。最自然的线程映射一眼就能看出来：给输出元素 $C_{ij}$ 分配一个线程，让它读 $A$ 的第 $i$ 行、 $B$ 的第 $j$ 列，做点积，再加上旧的 $C_{ij}$ ，写回结果。

对应的最小内核长这样：

cpp

__global__ void sgemm_naive(
    int M, int N, int K,
    float alpha,
    const float* A,
    const float* B,
    float beta,
    float* C) {

    const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < M && y < N) {
        float tmp = 0.0f;
        for (int k = 0; k < K; ++k) {
            tmp += A[x * K + k] * B[k * N + y];
        }
        C[x * N + y] = alpha * tmp + beta * C[x * N + y];
    }
}

这段代码看起来完全合理。它在数学上忠实，容易理解；因为输出矩阵是二维的，所以用二维网格和二维线程块；每个线程职责清楚；内层循环就是教科书式的点积。没有任何功能性错误——而这恰恰是它如此有教育意义的原因：代码正确、工作负载理论上极好、性能却仍然糟糕。 问题不在正确性，在映射。

这是 CUDA 开始区别于普通算法设计的地方。算法课上，一份功能正确的矩阵乘法已经算主要成果；高性能 GPU 编程里，功能正确只是起点。真正的问题是：这个映射产生的那条数据流，是否尊重这台机器？

3. 理论上，矩阵乘法应当是计算受限的

课程接着问了一个看似简单、实则极其尖锐的问题：矩阵乘法到底是计算受限（compute-bound）还是带宽受限（bandwidth-bound）？

在数学层面，矩阵乘法看起来像是 GPU 的梦想负载。算术量按 $O (N^{3})$ 增长，而存储的数据量只按 $O (N^{2})$ 增长。矩阵越大，算术强度越高。幻灯片把这个论断落到实处，用的是 H100 SXM 和 $4096 \times 4096$ 的矩阵乘法：H100 有约 80 GB HBM3，峰值内存带宽约 $4.096 TB/s$ ，FP32 CUDA 核吞吐约 $60 TFLOP/s$ （低精度下张量核还能更高）。

对 $C = A \times B + C$ ，其中 $A, B, C \in R^{4096 \times 4096}$ ，总浮点工作量约为

2 \cdot 409 6^{3} + 409 6^{2} \approx 137 GFLOPs,

系数 2 来自点积中每步一次乘法加一次加法，末尾的 $409 6^{2}$ 计入把旧 $C$ 加进来的那步。幻灯片明确给出了这个数字。

至少要被读进来的数据是三个输入矩阵 $A$ 、 $B$ 和旧的 $C$ ，合计约

3 \cdot 409 6^{2} \cdot 4 B \approx 201 MB,

写回输出 $C$ 还要

409 6^{2} \cdot 4 B \approx 67 MB .

所以总的下界流量约 $268 MB$ ——这组数字同样是幻灯片原样给出的。

往下算理想下界就很容易了。 $268 MB$ 除以 $4.096 TB/s$ ，传输下界约 $0.07 ms$ ； $137 GFLOPs$ 除以 $60 TFLOP/s$ ，计算下界约 $2.3 ms$ 。也就是说，纸面计算说计算时间大约是纯传输下界的三十倍。课程因此在算法层面给出正确结论：GEMM 应当是计算受限的，而不是带宽受限的。

这里是把课程接上 roofline 模型的最佳位置。Roofline 说性能被

P \leq min (P_{peak}, B W_{peak} \cdot I)

封顶，其中 $P_{peak}$ 是峰值算术吞吐， $B W_{peak}$ 是峰值带宽， $I$ 是算术强度（每字节 FLOPs）。

代入具体数字可以让这个上界更直观。GEMM 的算术强度

I = \frac{137 GFLOPs}{268 MB} \approx 512 FLOP/B,

把它乘上 H100 的峰值带宽 $4.096 TB/s$ ，得到 $B W_{peak} \cdot I \approx 2.1 PFLOP/s$ ——比 $60 TFLOP/s$ 的 FP32 峰值高出一个数量级以上。也就是说，只要数据搬运模式接近理想，内存屋顶对应的那条斜线就远在计算屋顶之上，GEMM 理应贴着计算屋顶跑，而不是挂在内存屋顶（memory roof）下。这个 512 FLOP/B 的量级，也正是 GEMM 在所有主流工作负载里算术强度名列前茅的原因。

然而朴素内核并没有这么做。 这就把我们带进了真实世界。

4. 欢迎来到真实世界：两个数量级的差距

幻灯片上赫然写着「Welcome to the real world!」。实测结果是：H100 的 FP32 CUDA 核理论上有约 $60 TFLOP/s$ ，朴素内核却只跑到约 $500 GFLOP/s$ ， $4096 \times 4096$ 那次运行总耗时超过 200 毫秒。

这不是一点小低效，这是性能崩塌。实测吞吐比 FP32 峰值低了两个数量级以上。课程毫不客气地把它当成一个调试谜题：算法正确，算术负载很大且高度并行，纸面计算预测它是计算受限的——那为什么实现表现得这么糟？

把时间对一对也很扎眼：纸面下界是 $2.3 ms$ ，实测超过 $200 ms$ ——差了将近一百倍，而这两个数字对应的算法是同一个，问题只出在实现方式上。这本身就是一记警钟：在 GPU 上，「算法」和「算法的高性能实现」是两个物种。

答案微妙但根基性：纸面计算数的是「逻辑」数据移动，而不是实现真正诱导出的「物理」数据移动。 下界流量估计隐含两个假设：内存交付的每个字节都有用，该重用的数据确实被重用了。朴素内核同时违背了这两条——它产生了浪费带宽的糟糕访问模式，又反复重载本应被共享的值。一旦如此，实现的算术强度就塌了，一个理论上计算受限的工作负载，会表现得像一个低效的内存内核。

这就是为什么课程没有直接跳去用「共享内存」这个魔法招数，而是先问一个更基本的问题：我们真的理解内存吗？

5. DRAM、HBM、行缓冲区，以及合并访问为什么存在

讲解稿花了大量篇幅拆解 DRAM 的物理现实，这一步走对了。跳过它，合并访问（coalesced access）听起来就像一条要死记的 CUDA 规则；理解了 DRAM，合并访问就变成一件必然发生的事。

直觉。 在电路层面，一个 DRAM 单元就是一个电容加一个晶体管。电容存着一丁点电荷；由字线（wordline）控制的晶体管把单元接到位线（bitline）上。读一个单元，就是把微小的单元电容接到大得多的位线上，再去感应由此产生的电压扰动。因为存的电荷太少，信号很弱，必须放大；更糟的是这种读本质上是破坏性的，感应之后还得把原值写回去。

于是下一个架构事实顺理成章：DRAM 不是为逐个读孤立的标量字优化的，它是围绕「行」优化的。 一次访问先激活一行、把它搬进感应放大器，再从那个临时行缓冲（row buffer）里服务突发传输（burst transfer）。一行打开之后，邻近访问就很便宜；打开另一行则很贵——系统必须预充电（precharge）、激活、感应，事后还要恢复新行。

把这段话再翻译成工程直觉，就是：凑够一整行的连续请求，比你发一百个零散请求要便宜得多。 这跟火车运输是同一门账——一趟专列把一整条车厢的货物拉到目的地，比一百趟单件快递划算，因为发车的固定成本太高。DRAM 里「发车」的成本就是行激活。

关键约束。 DRAM 里昂贵的操作不是「取一个 float」，而是「激活一个新行」。所以内存系统希望每次行激活都尽可能服务更多有用的邻近请求。HBM 一节把这说得更具体：HBM 把几十个 DRAM die 垂直堆叠、用硅通孔（TSV）连通，靠大量 bank 并行、更短的走线和更宽的接口把带宽推到 TB/s 量级——但它在本质上仍是 DRAM：依旧按行组织，依旧有行缓冲。一个 bank 有一个约 1–2 KB 的行缓冲或页面，突发传输从活动行里成块返回数据；channel、stack、众多 bank 并行工作，才堆出 HBM 的总带宽。但所有这些令人惊叹的带宽，都依赖软件给出能让行缓冲和突发传输被高效利用的地址。

这就是步幅（stride）为什么这么重要。幻灯片里有一个带宽实验：小步幅约 $1418 GB/s$ ，病态步幅——相当于打开一整页却只消耗其中一个字——掉到约 $111 GB/s$ ，在该测试里只有峰值的 8%。两者相差约 12–13 倍。

到这一步，合并访问的真正含义就清楚了：合并访问不是编程风格偏好，而是软件侧的条件——它让 DRAM/HBM 用少数几次密集的、来自已打开行的突发事务，就满足一个 warp 的请求。 warp 的请求连续，取回的每个块大多是有用数据；warp 的请求跨步、离散、散落，硬件就反复开行、搬运内核几乎不用的数据。性能就是这样崩的。

由此得到本章第一个诊断视角：别问「一个线程在访问什么」，要问「一个 warp 的 32 条通道一起在请求什么地址」。

6. warp、锁步执行，以及为什么线程局部性不够

下一层关键机制是 warp 执行模型。CUDA 程序员常写二维、三维线程块，因为问题本身是二维、三维的；但硬件并不直接执行「一个 2D 块」，它执行 32 线程的 warp，且通道按线性顺序分配：先 threadIdx.x，再 threadIdx.y，最后 threadIdx.z。warp 组织那页幻灯片强调的就是这条规则。

这很重要，因为内存系统看到的不是某一个线程的访问轨迹，而是 warp 里 32 条通道锁步发出的联合访问模式。单个线程的访问模式可能看起来完全合理，但它所属的 warp 的集体访问模式可能是灾难——朴素矩阵乘法的映射正是如此。

只盯着计算 $C_{ij}$ 的单个线程，它的访问看起来混杂但可以理解：它沿 $A$ 的第 $i$ 行前进（在行主序里是连续的），沿 $B$ 的第 $j$ 列下行（跨步，不局部）。单线程视角下， $A$ 很好， $B$ 很差。

但硬件问的不是这个问题。

假设 warp 里相邻的通道对应输出矩阵 $C$ 的相邻行，列 $j$ 相同。那么在 $k$ 循环的某一步，32 条通道读的是：

$A [i, k], A [i + 1, k], A [i + 2, k], \dots$ ——这是行主序存储下 $A$ 的一列，跨步极大；
以及同一个 $B [k, j]$ ，这往往可以由类似广播的行为服务。

这就是自然映射里的深层 bug。线程局部推理说：「 $A$ 是按行的， $A$ 应该没问题。」warp 级推理说：「在固定的内层迭代上，warp 其实是在沿 $A$ 的一列下行。」这彻底摧毁了全局内存局部性。幻灯片序列——「单个线程」「一个 warp 的线程（这为什么糟？）」「这就是我们得到的」——画的正是这个错配。

停下来想想

亲手验证一下。拿一张方格纸，行主序画出 $A$ ；再想象 32 个编号相邻的线程横排开（相邻 threadIdx.x），它们各自负责 $C$ 里上下相邻的 32 行、同一列。内层循环第 $k$ 步，它们取的是 $A [i + k ?]$ … 自己标出这 32 个地址，你会看到它们在内存里均匀散开、彼此相隔 N 个 float。这就是「每个线程都合理，32 个线程一起就不合理」的物理演示。

这节最要紧的概念结论是：线程局部性是不够的。对全局内存而言，重要的是 warp 局部性。

7. 第一个真正的优化：让 warp 与行对齐

一旦问题在 warp 层面被说清，修复办法简单得惊人：不让 warp 覆盖输出 $C$ 的一列条带，而让它覆盖一行条带。

再考虑一次内层循环的某一步。所有通道用同一个 $k$ 。如果它们共享同一输出行 $i$ ，覆盖相邻的输出列 $j, j + 1, \dots, j + 31$ ，那么：

所有通道需要同一个标量 $A [i, k]$ ，天然适合广播；
所有通道需要连续的 $B [k, j], B [k, j + 1], \dots, B [k, j + 31]$ ，正好是 $B$ 第 $k$ 行的连续片段。

这正是内存系统想要的：一个操作数变成广播式，另一个变成合并式。幻灯片「一个 warp 的线程（我们想要的）」展示的正是这种改进后的 warp 级访问模式。

这个优化最美妙的地方在于，相对它的效果，代码改动小得几乎可笑。算术没变，循环结构没变，算法没变——变的只是从线程坐标到输出坐标的映射。实测性能却从约 $0.5 TFLOP/s$ 跳到约 $6.3 TFLOP/s$ ，加速倍数和病态步幅实验里观察到的 12–13 倍带宽损失几乎一模一样。幻灯片把这个结果明确写出来，并总结为 Lesson 1：全局内存里的数据应该连续访问，以获取最大带宽效率。

这是全章第一个教学回报。一次看似微小的索引改动带来巨大加速，是因为这个改动编码了一条关于 HBM 和 warp 执行的大物理事实。CUDA 优化常常给人「代码小、效果大」的感觉，原因正在于此。

值得把这次改动的物理含义再放大看一眼：改动前，一个 warp 的每次访问都会让 HBM 打开一整页、却只消费其中一两个字；改动后，同一个 warp 取回的一整块突发数据几乎全部被用掉。带宽没有被「变多」，是被「用对了」。这也顺带解释了为什么加速倍数（12.6 倍）和病态步幅实验里的带宽损失（12–13 倍）几乎一样——你修掉的正是同一份浪费。

顺带补一个有用的解释点：GEMM 的算法算术强度没变，变的是 HBM 看到的实现算术强度。每条 HBM 事务里有用数据的占比提高后，内核不再浪费带宽，开始向纸面模型预测的计算屋顶移动。

8. 下一个瓶颈：HBM 变好了，但数据仍然没有被共享

12 倍加速很惊人，但离 H100 的峰值还很远。下一个问题很自然：现在瓶颈在哪？

幻灯片的回答是一个简单观察：线程取来的数据没有被共享。

这是下一个重要的系统洞见。即使修好了全局内存合并访问，内核仍然是让每个线程把需要的数据从 HBM 取进自己私有的寄存器。寄存器是私有的。相邻线程如果需要 $A$ 和 $B$ 的重叠片段，这份重叠不会被利用——相同的值会被不同线程各自加载很多次。

算法层面，GEMM 应该搬运 $O (N^{2})$ 的数据、做 $O (N^{3})$ 的算术；但如果每个线程都私装一份重叠操作数，实现的全局内存流量就会变得接近 $O (N^{3})$ 。这正是开头那页幻灯片警告过的通信放置失败：算术没变，重用的放置却完全错了。

这也是课程引入共享内存的理由。

9. 共享内存平铺：用块级协作减少 HBM 流量

共享内存是容量小、速度快、块内所有线程可见的片上存储。关键词是共享：寄存器是私有的，撑不起线程间的重用；共享内存可以。

直觉。 平铺（tiling）的思路朴素直白。假设一个线程块负责输出矩阵 $C$ 的一块 tile。要算这块输出，块内所有线程需要 $A$ 里对应的一块 tile 和 $B$ 里对应的一块 tile。与其让每个线程都独立从 HBM 取自己那份，不如让线程块协作把输入 tile 搬进共享内存，然后每个线程从共享内存反复重用这些值来累积部分结果。

幻灯片把过程拆成了显式步骤：确定网格、块、线程的排布；确定这个块拥有 $C$ 的哪块输出 tile；确定 $A$ 、 $B$ 里对应的 tile 区域；把 tile 从全局内存拷进共享内存；逐块计算部分结果；沿 $K$ 维重复；最后把结果 tile 写回 $C$ 。简化后的内核骨架如下：

cpp

__shared__ float As[BM][BK];
__shared__ float Bs[BK][BN];

float acc = 0.0f;

for (int tile = 0; tile < K; tile += BK) {
    As[ty][tx] = A[row * K + (tile + tx)];
    Bs[ty][tx] = B[(tile + ty) * N + col];
    __syncthreads();

    for (int k = 0; k < BK; ++k) {
        acc += As[ty][k] * Bs[k][tx];
    }
    __syncthreads();
}

具体索引取决于 tile 尺寸和块布局，但骨架相同：每个 tile 从 HBM 到共享内存搬一次，之后大量算术操作在片内重用这份 tile。tile 尺寸 $B M, B K, B N$ 本身也是要调的参数：太小，块内重用不够，平铺带来的收益有限；太大，共享内存装不下，还会压掉占用率。经典的选择如 $16 \times 16$ 、 $32 \times 32$ 、 $32 \times 8$ ，都是带宽、容量、占用率之间的平衡点——这个问题到寄存器平铺后会再遇到一次，只是换了一个层级。

通信图景因此彻底改变。每个从 HBM 进共享内存的值，现在被块内多个线程使用。 $A$ 的一块 tile 被许多输出列重用， $B$ 的一块 tile 被许多输出行重用。用 roofline 的语言说，HBM 看到的算术强度提高了——同样多的 HBM 字节，现在支撑着更多的 FMA。精确一点：假设块内每个元素被复用 $B M$ （或 $B N$ ）次，HBM 流量就从「每线程读一份」降到「每块读一份」，降幅正是 tile 的复用度。

课程报告，共享内存平铺把性能从约 $6.3 TFLOP/s$ 推到约 $9 TFLOP/s$ ，配套幻灯片称之为「算术强度改善了」。Lesson 2 一句话说清核心：平铺利用每个 SM 的共享内存，促成了块内重用。

这背后还有一个更深的机器学习系统教训。共享内存不像 L1 缓存，它是由软件管理的。也就是说，程序员或编译器要负责决定「什么该住在那里」。这是一种显式的数据编排（data orchestration）。当工作负载变复杂，哪些重用该用软件管理的存储来表达、哪些可以交给硬件缓存策略——这本身就成了一个重要的系统设计问题。

10. 共享内存比 HBM 快，但它也不是免费的

加上共享内存平铺之后，瓶颈再次转移。这是分层性能调试（hierarchical performance debugging）的绝佳案例：优化层次结构里的一层，会暴露出下一层。

性能分析幻灯片显示，内核仍然是内存受限的，但压力从 HBM 转移到了共享内存。warp stall 那页幻灯片解释：warp 停在那里，是因为 MIO（memory input/output，内存输入输出）指令队列满了，起因是块内密集的共享内存指令流量。「访问很多——仍然内存受限」这组幻灯片强调的正是这个新诊断。

为什么会这样？因为在基础平铺内核里，每个线程仍然只算一个输出元素。内层循环每走一步，它都要从共享内存取一项 $A$ tile 的值、一项 $B$ tile 的值进寄存器，做一次乘加，然后重复。共享内存比 HBM 快得多，但如果内核产生海量共享内存流量，共享内存管线自己就变成了限制器。

这是全章第二个性能调试教训：瓶颈没有消失，只是向层次结构内部移动了。

接下来的修复不是进一步减少 HBM 流量，而是减少每单位算术对应的共享内存流量——这把我们带到寄存器平铺。

11. 寄存器平铺：让每个线程计算不止一个输出

课程把下一个想法框得很好：想要更少的共享内存指令，每个线程就得做更多的事。

直觉。 这就是寄存器平铺（register tiling）背后的基本想法。现在的平铺内核里，一个线程只算一个输出 $C_{ij}$ ，从共享内存取来的值用一次就扔。但寄存器是整台机器上最快的存储。如果每个线程负责多个输出，一个从共享内存载入寄存器的值，在被丢弃前就能被多次 FMA 重用。每 FLOP 的共享内存流量随之下降。

11.1 一维寄存器平铺

一维寄存器平铺里，每个线程算一短串输出，而不是一个标量。例如，与其算某列里的一个输出元素，不如算几个竖直相邻或水平相邻的输出（取决于块和 tile 怎么组织）。线程在一个小寄存器数组里持有多个部分和：

cpp

float threadResults[TM] = {0.0f};

于是线程载入一个操作数片段后，可以在多个累加器之间重用。共享内存 tile 被进一步拆成更小的片段，变成线程私有的寄存器 tile；线程算多个输出，而不是只算一个。

性能影响很大。课程报告，一维寄存器平铺把性能从约 $9 TFLOP/s$ 推到接近 $20 TFLOP/s$ ，又翻了一倍多。配套幻灯片说：共享内存带宽的利用被显著改善了。

11.2 二维寄存器平铺

一维寄存器平铺只是第一步。二维寄存器平铺让每个线程算一小片输出，而不是一维向量。这意味着 $A$ 侧和 $B$ 侧的操作数都能在线程内部被更有效地重用。

到这一步，每个线程变成了一个微型矩阵乘法微内核：从 $A$ 载入一小片片段，从 $B$ 载入一小片片段，在寄存器里的 $T_{M} \times T_{N}$ 累加器 tile 上做一个小的外积更新。这已经非常接近优化库中高性能 GEMM 内核的结构了。

课程报告，二维寄存器平铺把内核再推一步，到约 $25.8 TFLOP/s$ 。幻灯片明确说，二维平铺进一步改善了共享内存带宽利用。

为什么二维比一维又进一大步？一维平铺下，线程持有一串输出（比如 $T_{M}$ 个竖直相邻的输出），载入 $A$ 的一个片段可以横跨这串输出重用，但 $B$ 的每个元素仍然只被用一次。二维平铺把重用的受益面扩大到两边： $A$ 片段被 $T_{N}$ 个累加器共用， $B$ 片段被 $T_{M}$ 个累加器共用，每个从共享内存取出的值都贡献 $T_{M} \cdot T_{N}$ 次 FMA。这正是「每 FLOP 的共享内存流量再降一个数量级」的来源。回想第 2 章说过的向量化思想——把有用工作与簿记的比值拉高——在这里重演了：只是这次的「簿记」是共享内存指令。

11.3 为什么寄存器平铺有效

寄存器平铺之所以有效，是因为它把重用又向层次结构深处推了一层：

合并访问让 HBM 访问变高效；
共享内存平铺让 HBM 数据在块内可重用；
寄存器平铺让共享内存数据在线程内可重用。

通信量不只是抽象意义上变少了，而是被移到了越来越便宜的层次。

当然这里也有取舍：寄存器平铺提高寄存器压力（register pressure）。每个线程干的活越多，占用的寄存器越多，占用率（occupancy）可能下降——因为一个 SM 上能同时驻留的 warp 变少了。这也是为什么最优 GEMM tile 尺寸从来不是固定的「神奇数字」：它取决于机器、精度格式、可用共享内存和目标占用率区间。「最好」的配置永远是重用、资源占用、调度灵活性三者之间的折中。

从这一刻起，CUDA 编程开始变得非常像系统设计，而不是简单的并行化。

12. 访存向量化：让每条指令搬更多字节

合并访问、共享内存平铺、寄存器平铺之后，课程简要介绍了访存向量化（memory access vectorization），这是又一个有用的细化。

即使访问已经合并，一次只载一个 float 仍然会产生超过必要数量的内存指令。如果对齐条件允许，可以载更宽的块，比如 float2、float4，减少指令数、提高每条指令的有效载荷。

课程还点出一个重要的布局技巧：一个操作数天然按行组织，另一个天然按列组织。在从全局内存拷进共享内存的过程中，可以旋转或转置暂存布局，让之后从共享内存进寄存器的访问变成按行、更局部、更容易向量化。这是一个很重要的系统洞见：共享内存不只是全局内存的加速拷贝，它还可以当布局转换缓冲区（layout transformation buffer）用。 软件用 HBM→SMEM 的这次搬移，不只把数据搬进来，还把数据重排成对管线下一阶段更友好的形状。

这是看待优化内核里片上内存最有价值、却最不显眼的角度之一：它们不只是用来存储的，它们把数据重排成下一层层次结构喜欢的形态。

13. warp 平铺：让 warp 本身成为一个有意义的计算对象

课程接着引入 warp 平铺（warp tiling），这是又一次概念升级。

到目前为止，讨论一直围着块和线程打转。但硬件真正调度的是 warp。所以一个成熟的映射，不该只给块和线程分派有用的工作，也应该把有用的工作分派给warp 这个集体对象。

warp 平铺的定义很明确：一个 warp 协作计算输出矩阵的一个子 tile，而不是让每个线程独立算一个元素。它的目的包括：让计算组织与 warp 执行单元对齐、减少不必要的同步、更有效地利用寄存器重用、进一步抬高算术强度。

这是一个漂亮的概念收束，因为它接回了早先的「执行错配」：硬件执行的是 warp，不是任意的二维线程抽象。因此好的软件应该越来越显式地围绕 warp 这个一等实体来组织计算。

warp 平铺天然位于块级共享内存平铺和线程级寄存器平铺之间。块在共享内存里拥有一块大 tile；块内每个 warp 拥有更小的子 tile；warp 内每个线程再拥有更小的寄存器 tile。这个三层分解，正好镜像了机器真实的执行与存储结构。

但课程也提醒，warp 平铺会引入新麻烦：一旦 warp 以更结构化的方式访问共享内存，这些访问模式可能触发存储体冲突（bank conflict）。到这儿，「魔鬼在细节里」这句话就完全应景了。

14. 共享内存存储体冲突：为什么「快内存」也可能变慢

共享内存快，是因为它分 bank。H100 上，共享内存实际上组织成 32 个 bank。理想情况下，一个 warp 的 32 条通道各访问一个不同的 bank，整次 warp 载入一拍完成；若多条通道落在同一个 bank 上，访问串行化；最坏情况下 32 条通道争同一个 bank，本该一拍完成的访问要多花很多拍。

H100 共享内存那页幻灯片把意思总结得很清楚：无冲突访问一拍，广播也一拍，但 $k$ 路冲突要花 $k$ 拍——所以性能取决于 bank 映射效率，而不仅仅是内存大小。

这引入了新一类局部性问题。即使数据在共享内存里「彼此靠近」，只要它们和 bank 映射错位得厉害，访问照样慢。这常发生在数据按一种朝向写入、按另一种朝向读出时：一个 warp 按行写一块 tile，之后却按列读它；步幅一旦和 bank 映射交互得很糟，访问模式就高度冲突。

停下来想想

自己推一遍最经典的那个冲突。共享内存按 地址 % 32 分到 32 个 bank。假设一个 $32 \times 32$ 的 tile，每行 32 个 float，按行存放。一个 warp 去读第 $k$ 列的全部 32 行——线程 $i$ 读地址 $i \cdot 32 + k$ 。取模后，所有 32 个地址都落在同一个 bank $k$ 上：32 路冲突，一拍变成 32 拍。这就是「按行写、按列读」的病态。解决方案就是给它换个朝向：要么每行补 1 个填充元素，把步幅改成 33；要么干脆在从 HBM 拷进共享内存时就把转置做掉（第 12 节的布局变换），之后按行读就天然无冲突。两种解法的本质是同一个——让 地址 % 32 重新散开。

经典解药是填充（padding）。如果共享内存 tile 宽度正好是 32，warp 又按列读，所有通道可能撞上同一个 bank。每行多补一个填充元素，步幅从 32 变成 33，病态对齐就被打破。另一个解药是借向量化那节提到的思路，在暂存时改布局，让之后的访问天然对 bank 友好。

课程的结论很准：共享内存的性能不只取决于用了多少，还取决于 warp 如何映射到它的物理 bank 上。 一旦内核优化到一定程度，这些细粒度细节就开始承重。

15. 优化阶梯的统一视角

走到这里，课程已经介绍了一长串具体技术：合并访问、共享内存平铺、寄存器平铺、向量化访问、warp 平铺、规避存储体冲突。把它们当孤立的技巧去背，会错过更深的结构。

这些技术形成了一条连贯的优化阶梯：

全局内存合并访问修复行主序布局与 warp 通道顺序之间的错配；
共享内存平铺修复算法重用与线程私有寄存器之间的错配；
寄存器平铺修复块级重用与线程算术吞吐之间的错配；
向量化修复标量指令粒度与内存系统偏好的传输宽度之间的错配；
warp 平铺修复逻辑块分解与真实执行单元（warp）之间的错配；
存储体冲突管理修复逻辑共享内存局部性与片上 SRAM 物理 bank 结构之间的错配。

从这个角度看，本章的每一项优化都在做同一件事：消除工作负载的自然结构与硬件的实际结构之间的一处具体错配。

所以开头那页框架如此重要。CUDA 编程不是一袋技巧，它是把高维工作负载映射到线性、分层硬件上的纪律性过程——每一个性能问题，都能追溯到映射里的一处错配。

16. 为什么我们仍然打不过 cuBLAS

到课程结尾，优化过的 CUDA 核 GEMM 达到约 $25$ – $26 TFLOP/s$ ，比起步的 $0.5 TFLOP/s$ 提升了近两个数量级。幻灯片把这明确框定为把所有环节拼在一起的结果：连续全局内存访问、共享内存重用、warp 级与寄存器级平铺、向量化、bank 感知的共享内存访问。

但即便如此，它仍然没有完全摸到 H100 的理论 CUDA 核 FP32 峰值，更不用说张量核路径能到达的高度。幻灯片半开玩笑地把它总结为 Lesson 4：想在 cuBLAS 自己的主场上赢它，很难。

为什么？原因有好几条。

其一，生产库用的不是新想法，而是极端成熟和调优深度的旧想法。 cuBLAS 这类库会针对每种架构、每个问题规模仔细挑选 tile 尺寸——同样的思想，它们打磨到了手工难以复现的深度。这与 CUTLASS（NVIDIA 的开源 CUDA GEMM 模板库）文档里呈现的结构一致：warp 级 tile、寄存器分块、指令选择都是经过大量实测的。

其二，它们用软件流水线（software pipelining）和双缓冲（double buffering）。 加载下一块 tile 与计算当前 tile 重叠；在新架构上，cp.async 这类异步拷贝指令让重叠更彻底。CUTLASS 文档把这种「算一块、载一块」的流水线当作第一公民特性。

其三，它们经常直接打张量核（tensor core）。 张量核走不同的指令格式（例如 mma 指令），只要数据类型和 tile 形状对得上硬件，峰值吞吐就远高于 CUDA 核。CUTLASS 就是围绕 mma 指令族构建的。这一步的差距是代差级的：H100 上 FP32 的 $60 TFLOP/s$ 只是 CUDA 核的吞吐，张量核在混合精度下可以把这个数字翻一个数量级以上——我们的手工内核再怎么打磨 CUDA 核，天花板也还够不着那条路径的起跑线。

其四，它们还榨取了机器特定的调度、寄存器阻塞、指令选择和布局变换。 这些细节的深度，手工很难在合理时间内完整复现。

所以对这门课来说，正确的解读不是沮丧于「我们的内核还不是 cuBLAS」，而是尊重「cuBLAS 为什么难打」。当你亲眼看到必须同时对齐多少层映射与重用，自动调优、编译器支持和厂商优化库的必要性就再明显不过了。

旁证：FlashAttention

同样的逻辑还能走出 GEMM。算子融合（operator fusion）的目标，是让中间结果不落回 HBM。FlashAttention（Dao et al., 2022）把 $Q K^{T}$ 、softmax、 $O = S V$ 融进一个内核，只把最终输出写回 HBM，避开中间注意力矩阵的读写往返——这正是「数据尽量留在片上」这一章主旨的注意力版本。

讲解稿也正是在这里指向下一讲——机器学习编译器。一旦你理解了内核调优的复杂性，一个自然的问题就浮出来：这条优化阶梯里，至少有一部分能不能被自动生成或自动搜索出来？

结语：CUDA 编程，是尊重机器的映射

现在，可以把全章最深的教训用更成书面的方式重述一遍。

第一，理论算术强度不等于实现算术强度。 GEMM 在数学上是计算受限的，但一次糟糕的映射能让实现表现得像一个蹩脚的内存内核。

第二，优化的单位不是单个线程。 真正的基本单位是分层的：HBM 页、内存事务、warp、线程块、共享内存 tile、寄存器 tile、执行管线。优化 CUDA，意味着理解这全部。

第三，重用必须不断向内搬。 在 HBM 里重用太贵，共享内存里重用更好，寄存器里重用最好。高性能内核，就是把重用逐步推向层次结构更深处。

第四，工作负载的几何不等于执行的几何。 矩阵是二维的，warp 是一维的。任何无视这一点的线程映射，都很可能一开始就浪费带宽。

第五，很小的代码改动可以编码很大的硬件事实。 第一次 12 倍加速来自一次微小的索引重映射，因为那次重映射从根本上改变了 warp 让内存系统去做的事。

第六，性能调优是迭代式的瓶颈调试。 先出问题的是 HBM，然后是共享内存，然后是共享内存指令压力，再然后是 bank 映射和传输粒度。没有一劳永逸的招数，只有一条不断被揭示的瓶颈序列。

这也正是 CUDA 编程应该出现在机器学习系统课程里的原因。它教会你：性能不是算法单方面的属性，而是算法在一个受硬件约束的执行环境里被实现出来之后的属性。

课程的中心句，现在可以在全部深度上被理解了：

CUDA 编程，就是把高维工作负载映射到线性、分层硬件之上的问题。

这不是一句口号，它是本章每一项优化背后的组织原则。合并访问，是 warp 尊重 HBM 的突发式组织；共享内存平铺，是块级重用机会被放到片上；寄存器平铺，是人们意识到共享内存本身也能变成瓶颈；warp 平铺，是认真对待硬件执行单元；规避存储体冲突，是尊重那个「快」内存自身的内部结构。

这也是本章作为编译器话题桥梁的原因。一旦你看清有多少细节在承重——布局、warp 映射、平铺、同步、向量化、bank 映射、寄存器压力、占用率——内核生成和自动调优就显得不是奢侈品，而是近乎必然。这个搜索空间太丰富，光靠直觉在大规模上根本走不完。

所以正确的最终心智模型不是「CUDA 让我启动很多线程」，而是：

CUDA 让我设计数学结构与机器结构之间的映射。

映射越尊重机器，性能就越从理论靠近现实。

第 5 章 CUDA 编程作为软硬件协同优化 ​

从朴素矩阵乘法到分层分块 ​

1. CUDA 编程首先是映射问题 ​

2. 运行示例：一个最朴素的矩阵乘法内核 ​

3. 理论上，矩阵乘法应当是计算受限的 ​

4. 欢迎来到真实世界：两个数量级的差距 ​

5. DRAM、HBM、行缓冲区，以及合并访问为什么存在 ​

6. warp、锁步执行，以及为什么线程局部性不够 ​

7. 第一个真正的优化：让 warp 与行对齐 ​

8. 下一个瓶颈：HBM 变好了，但数据仍然没有被共享 ​

9. 共享内存平铺：用块级协作减少 HBM 流量 ​

10. 共享内存比 HBM 快，但它也不是免费的 ​

11. 寄存器平铺：让每个线程计算不止一个输出 ​

11.1 一维寄存器平铺 ​

11.2 二维寄存器平铺 ​

11.3 为什么寄存器平铺有效 ​

12. 访存向量化：让每条指令搬更多字节 ​

13. warp 平铺：让 warp 本身成为一个有意义的计算对象 ​

14. 共享内存存储体冲突：为什么「快内存」也可能变慢 ​

15. 优化阶梯的统一视角 ​

16. 为什么我们仍然打不过 cuBLAS ​

结语：CUDA 编程，是尊重机器的映射 ​

延伸阅读 ​