PMPP 第五章

30 Jun, 2026

Memory architechure and data locality

5.1 Memory bandwidth as a performance limiter

CUDA 程序的性能不一定总是被计算单元限制，很多时候真正的瓶颈是 global memory 访问。也就是说，GPU 可能不是“算不动”，而是“数据供不上”。这一节主要建立一个性能分析框架：通过 compute-to-global-memory-access ratio 判断 kernel 更可能是 compute-bound 还是 memory-bound。

1. GPU 性能的两个常见约束

GPU 在单位时间内能完成的工作有硬件上限，最常见的是两个指标。

峰值计算吞吐量表示 GPU 每秒最多能执行多少算术运算，通常用 FLOPS 表示。例如书中提到 H100 的单精度峰值计算吞吐量是：

66.9 TFLOPS

峰值内存带宽表示 GPU 每秒最多能从某种内存中读写多少字节数据，通常用 GB/s 或 TB/s 表示。例如 H100 的全局内存带宽是：

3.35 TB/s

简单理解：

计算吞吐量：GPU 每秒最多能算多少
内存带宽：GPU 每秒最多能搬多少数据

CUDA kernel 的性能到底受哪个限制，取决于这个 kernel 的计算量和访存量之间的比例。

2. Compute-bound、Memory-bound 与算术强度

如果一个 kernel 的性能主要受计算单元限制，称为 compute-bound。这类 kernel 中，计算单元大部分时间都在忙，内存能够持续提供数据，程序瓶颈在于“算得够不够快”。

如果一个 kernel 的性能主要受内存带宽限制，称为 memory-bound。这类 kernel 中，内存通道大部分时间都很忙，计算单元经常等待数据到达，程序瓶颈在于“数据能不能及时送到”。

判断这两种瓶颈的核心指标是：

compute-to-global-memory-access ratio = 浮点运算次数 / 全局内存访问字节数

单位是：

FLOP/B

这个指标也常被称为：

arithmetic intensity
computational intensity

也就是算术强度。

算术强度越高，说明每搬 1 字节数据能做更多计算，更容易成为 compute-bound；算术强度越低，说明搬很多数据只做很少计算，更容易成为 memory-bound。

区分 memory-bound 和 compute-bound 的大致阈值可以用硬件峰值计算吞吐量除以峰值内存带宽得到：

阈值 = 峰值计算吞吐量 / 峰值内存带宽

以 H100 为例：

66.9 TFLOPS / 3.35 TB/s = 20.0 FLOP/B

所以在 H100 上可以粗略认为：

算术强度 < 20.0 FLOP/B：更可能是 memory-bound
算术强度 > 20.0 FLOP/B：更可能是 compute-bound

不过这只是理论判断。实际性能还会受到访存合并、缓存命中、指令效率、occupancy、调度开销等因素影响。

3. Roofline Model 和 speed-of-light analysis

Roofline Model 是一种分析程序性能瓶颈的模型。它用一张图表示 kernel 的算术强度和实际计算吞吐量之间的关系。

横轴：算术强度，FLOP/B
纵轴：实际计算吞吐量，GFLOPS 或 TFLOPS

图中通常有两条“屋顶线”。

第一条是水平线，由硬件峰值计算吞吐量决定。第二条是从原点出发、具有正斜率的斜线，由硬件峰值内存带宽决定。两条线的交点就是 memory-bound 和 compute-bound 的分界点。

可以简单理解为：

左边：算术强度低，受内存带宽限制
右边：算术强度高，受计算吞吐量限制

在左侧 memory-bound 区域中，程序性能随着算术强度提高而提高，因为每搬一次数据能做更多计算；到了右侧 compute-bound 区域后，继续提高算术强度也无法突破计算峰值。

如果一个程序点离屋顶线很近，说明资源利用得比较充分；如果离屋顶线很远，说明还有优化空间。

书中还介绍了 speed-of-light analysis，也就是把程序实际达到的性能和硬件理论上限比较，看看离“理论极限”还有多远。

例如，两个长度为 10 亿的向量相加：

读取 A：4 GB
读取 B：4 GB
写入 C：4 GB
总访存量：12 GB

如果 kernel 执行时间是 4 ms，那么实际内存访问速率是：

12 GB / 4 ms = 3 TB/s

H100 的理论全局内存带宽是：

3.35 TB/s

所以带宽利用率是：

3 / 3.35 ≈ 90%

这说明该向量加法 kernel 已经接近硬件内存带宽上限，进一步优化空间不大。最快理论时间大约是：

12 GB / 3.35 TB/s ≈ 3.6 ms

所以如果实际已经跑到 4 ms，再想大幅优化就很难，因为它已经接近“带宽天花板”。

4. 向量加法：典型 memory-bound kernel

向量加法核心代码是：

C[i] = A[i] + B[i];

每个线程做的事情是：

读取 A[i]：4 B
读取 B[i]：4 B
写入 C[i]：4 B
执行一次浮点加法：1 FLOP

所以每个元素总共访问：

12 B

计算量是：

1 FLOP

因此算术强度为：

1 FLOP / 12 B = 0.083 FLOP/B

这个数远远小于 H100 的 20.0 FLOP/B，所以向量加法是非常典型的 memory-bound kernel。

也就是说，向量加法慢不是因为加法本身慢，而是因为每做一次加法就要搬很多数据。它的优化上限主要受 global memory bandwidth 限制。

5. 矩阵乘法：理论上 compute-bound，朴素实现却 memory-bound

矩阵乘法和向量加法不同，它有大量数据复用。对于矩阵乘法：

P = M × N

计算一个输出元素，需要做一组点积。每个输出元素大约需要：

N 次乘法 + N 次加法

也就是大约：

2N FLOP

如果是理想实现，每个输入元素只从全局内存读取一次，那么矩阵乘法的算术强度会很高。当矩阵规模较大时，它的理想算术强度可以远高于 H100 的 20 FLOP/B 阈值，因此矩阵乘法本来有潜力成为 compute-bound 计算。

简单说，矩阵乘法很适合 GPU，因为它具有这样的特点：

搬一次数据，可以反复计算很多次

但是，第三章中的朴素 CUDA 矩阵乘法实现没有利用好这种数据复用。核心循环类似：

for (int k = 0; k < Width; ++k) {
    Pvalue += M[row * Width + k] * N[k * Width + col];
}

每次循环中：

读取 M 的一个元素：4 B
读取 N 的一个元素：4 B
执行一次乘法 + 一次加法：2 FLOP

所以算术强度是：

2 FLOP / 8 B = 0.25 FLOP/B

这个数也远远小于 H100 的 20.0 FLOP/B。因此，虽然矩阵乘法理论上可以是 compute-bound，但这个朴素实现因为反复从 global memory 读取相同数据，实际变成了 memory-bound。

这一节最后引出的核心优化方向是：提高 CUDA kernel 性能，不能只关注线程数量和计算并行度，还要减少 global memory 访问。

尤其是矩阵乘法中，同一个 M 或 N 元素会被多个线程重复使用。如果每次都从 global memory 读取，就会浪费大量带宽。后面要学习的 shared memory 和 tiling，本质上就是为了：

把会重复使用的数据先搬到更快的片上存储中，
让多个线程复用它，
从而减少 global memory traffic。

本节核心可以概括为：

GPU 程序性能常常受内存带宽限制，而不是计算能力限制；
判断 compute-bound / memory-bound 的关键指标是算术强度 FLOP/B；
向量加法算术强度很低，是典型 memory-bound；
矩阵乘法理论上有高数据复用潜力；
朴素矩阵乘法没有利用数据复用，所以仍然 memory-bound；
shared memory 和 tiling 的目标就是减少 global memory 访问，提高算术强度。

5.2 CUDA memory types

这一节主要介绍 CUDA 设备中的几类内存：global memory、constant memory、local memory、registers、shared memory。不同内存类型的核心区别在于：访问速度、可见范围 scope、生命周期 lifetime，以及适合存放的数据类型。

1. CUDA 内存类型的整体关系

CUDA 设备中的内存大致可以分成两类：

片外内存：
global memory
constant memory
local memory 本质上也在 global memory 中

片上内存：
registers
shared memory

性能上可以粗略理解为：

registers 最快
shared memory 较快
global / local memory 较慢
constant memory 在合适访问模式下较快

其中，global memory 容量大，但位于片外 DRAM，访问延迟高、带宽有限，是很多 CUDA kernel 的主要性能瓶颈。前面章节中使用的输入输出数组，例如 M、N、P，通常都位于 global memory。

constant memory 主要用于存放只读数据。Host 可以写入，device 只能读取。它的变量必须声明在函数体外，生命周期贯穿整个 application，所有 grid 都可见。书中提到 constant memory 总大小限制为 65,536 B。它实际存储在 global memory 中，但会通过 cache 加速，在合适的访问模式下可以很快。

local memory 名字叫 local，但并不是片上高速内存。它是每个 thread 私有的存储空间，但实际位置在 global memory 中，所以访问延迟也接近 global memory。自动数组变量、寄存器溢出 spilled registers、线程调用栈中的部分数据，都可能放到 local memory 中。

2. Registers：线程私有的最快存储

registers 是 SM 内部的片上存储，访问速度最快，通常用于保存每个线程私有的标量变量。

比如 kernel 中的普通自动标量变量：

int row;
int col;
float Pvalue;

这些变量一般会放在寄存器中。每个线程都有自己的一份，其他线程无法访问。

寄存器快主要有几个原因。首先，它在芯片内部，访问延迟低、带宽高。其次，算术指令可以直接使用寄存器作为操作数，不需要额外的内存加载指令。例如：

fadd r1, r2, r3

如果操作数在 global memory 中，则需要先执行 load，再进行计算：

load r2, r4, offset
fadd r1, r2, r3

所以，把频繁访问的线程私有变量放在寄存器中，可以减少指令数量和访存开销。

不过，寄存器不是越多越好。每个 SM 的寄存器资源有限，如果每个线程使用太多寄存器，那么一个 SM 上能同时驻留的线程数就会减少，从而降低 occupancy。也就是说：

每线程寄存器用得少 -> SM 可以容纳更多线程
每线程寄存器用得多 -> SM 可以容纳更少线程

书中还对比了 CPU 和 GPU 的寄存器架构。CPU 线程切换时通常需要保存和恢复寄存器状态，而 GPU 为了支持 zero-overhead scheduling，会把 SM 上所有已调度线程的寄存器都保存在 register file 中。这样 warp 切换时不需要保存和恢复上下文，因此 GPU 的 register file 通常比 CPU 大得多。

3. Shared memory：block 内线程共享的高速存储

shared memory 也是片上内存，访问速度远快于 global memory，但通常慢于 registers。它最大的特点是：同一个 block 内的所有线程可以共享访问。

可以这样理解：

registers：每个线程自己的草稿纸
shared memory：一个 block 共享的黑板
global memory：所有线程都能访问的大仓库，但比较远、比较慢

shared memory 的典型用途是：让一个 block 内的线程协作，把 global memory 中会被重复使用的数据先搬到 shared memory 中，然后多个线程反复使用这份数据，从而减少 global memory 访问。

例如矩阵乘法中，一个 block 内多个线程会重复使用同一批 M 和 N 元素。如果每个线程都从 global memory 读取，就会产生大量重复访存；如果先把这些元素加载到 shared memory，再让 block 内线程复用，就可以明显减少 global memory traffic。

shared memory 的变量通过 __shared__ 声明：

__shared__ float tile[16][16];

它的特点是：

每个 block 有自己的一份
block 内所有 thread 可见
kernel 执行结束后失效
适合线程协作和数据复用

书中还提到 Hopper 架构开始支持 distributed shared memory。传统 shared memory 的作用范围是一个 block，而 distributed shared memory 允许同一个 thread block cluster 内的线程访问 cluster 中其他 block 的 shared memory，从而扩大线程协作范围。

4. 变量声明方式与内存位置

CUDA 中变量声明方式会影响变量被放到哪类内存中。

普通自动标量变量通常放在 registers 中：

int tx = threadIdx.x;
float Pvalue = 0.0f;

这类变量每个线程私有，生命周期是线程执行期间，访问速度很快。

自动数组变量可能放到 local memory 中：

float temp[16];

虽然它也是每个线程私有，但 local memory 实际在 global memory 中，所以访问可能较慢。不过，如果数组较小，并且访问下标是编译期常量，编译器可能会把它优化到 registers 中。

__shared__ 变量放在 shared memory 中：

__shared__ float tile[16][16];

它由一个 block 内所有线程共享，适合保存 block 内会被复用的数据。

__constant__ 变量放在 constant memory 中：

__constant__ float coeff[1024];

它必须声明在函数体外，kernel 内只读，所有 grid 可见，适合保存小规模只读输入数据。

__device__ 变量放在 global memory 中：

__device__ int counter;

它对所有 kernel 和 thread 可见，生命周期贯穿整个 application，但访问较慢。书中也提醒，global variable 容易降低代码模块化，并可能引入 bug，因此应谨慎使用。

5. 本节核心总结

5.2 节的核心是：CUDA 中不同内存类型有不同的速度、作用范围和生命周期，合理选择内存类型是优化 CUDA kernel 的基础。

变量声明	存储单元	范围	生命周期
除array外的自动变量	register	thread	grid
自动数组变量	local	thread	grid
__device__ __shared__ int Shared Var;	shared	block	grid
__device__ int globalVar	global	grid	application
_device\_ __constant__ int ConstVar	constant	grid	application

可以重点记住：

registers：
线程私有，最快，适合频繁访问的标量变量。

shared memory：
block 内共享，片上高速，适合线程协作和数据复用。

local memory：
线程私有，但实际在 global memory 中，访问慢，常由自动数组或寄存器溢出产生。

global memory：
容量大，所有线程可访问，但访问慢，是主要性能瓶颈来源。

constant memory：
只读，小容量，所有线程可见，适合小规模常量数据。

这一节为后面的 tiling 做铺垫：既然 global memory 慢，而 shared memory 快，那么可以把会重复使用的数据从 global memory 搬到 shared memory 中，让 block 内线程复用，从而减少 global memory 访问，提高 compute-to-global-memory-access ratio。

好的，这里按你当前进度理解为 5.3 节 Tiling for reduced memory traffic，继续整理读书笔记。小节不超过 5 个，并且全部使用三级标题。

5.3 Tiling for reduced memory traffic

1. Tiling 的基本思想：用小而快的内存缓存局部数据

这一节的核心是 tiling，也就是把大规模数据划分成较小的块，让每个小块能够放进 shared memory 中。这样做的原因是 CUDA 设备中存在一个基本矛盾：

global memory：容量大，但访问慢
shared memory：容量小，但访问快

所以优化思路就是：不要让线程反复从 global memory 读取同一批数据，而是先把一小块数据从 global memory 搬到 shared memory，再让 block 内线程复用它。

书中用类比说明 tile 的概念：一面大墙可以由很多小瓷砖覆盖；类似地，global memory 中的大数据也可以被划分成多个小 tile，每个 tile 能够放入 shared memory。

不过，并不是所有数据结构和 kernel 都适合 tiling。一个重要条件是：kernel 对这些 tile 的计算可以相对独立地进行。如果数据之间依赖关系很复杂，tile 之间不能独立处理，那么 tiling 就不一定适用。

2. 矩阵乘法中的重复访存问题

在朴素矩阵乘法中，每个线程负责计算输出矩阵 P 的一个元素。为了得到这个元素，线程需要读取矩阵 M 的一行和矩阵 N 的一列，然后做点积。

问题在于：同一个 block 内的多个线程会重复访问相同的 M 和 N 元素。

例如，一个 block 负责计算 P 中的一个小区域。这个区域里的多个输出元素会共享部分输入数据：

计算同一行输出的线程，会重复读取 M 的同一行元素
计算同一列输出的线程，会重复读取 N 的同一列元素

如果这些线程各自都从 global memory 中读取这些数据，就会造成大量重复访存。

如上图所示：在一个 2×2 的输出 tile 中，每个 M 和 N 元素都会被访问两次。如果让 block 内线程协作，把这些元素只从 global memory 读取一次，然后放入 shared memory 中复用，就可以把 global memory traffic 减少一半。

如果 tile 大小是 TILE_WIDTH × TILE_WIDTH，那么这种复用带来的 global memory 访问减少比例会更明显。书中指出，对于 32 × 32 的 tile，理论上可以把 global memory traffic 降到原来的 1/32。

3. Tiled matrix multiplication 的执行方式

tiled matrix multiplication 的基本流程是：一个 block 负责计算输出矩阵 P 的一个 tile，而每个 phase 加载一部分 M 和 N 的输入 tile 到 shared memory 中。

每个 phase 大致分为三步：

1. block 内线程协作，从 global memory 加载一块 M tile 和一块 N tile 到 shared memory
2. 每个线程使用 shared memory 中的数据，完成当前 phase 的部分点积计算
3. 进入下一个 phase，复用同一块 shared memory，加载下一组 M tile 和 N tile

也就是说，一个完整点积不再一次性从 global memory 中读取所有数据，而是被拆成多个阶段完成。

例如，对于矩阵乘法中的一个输出元素，原本需要遍历完整的 k 维度。使用 tiling 后，k 维度被切成多个 tile。每次只处理一个 tile 范围内的数据。处理完当前 tile 后，再进入下一个 tile。

在 Fig. 5.7 中，M 和 N 被划分成多个 input tile，这些 tile 由粗线标出。每个线程执行的 dot product 计算现在被划分成多个 phase。在每个 phase 中，一个 block 内的所有线程协作，把一个 M tile 和一个 N tile 加载到 shared memory 中。

这可以通过让 block 中的每个线程分别加载一个 M 元素和一个 N 元素到 shared memory 中来实现，如 Fig. 5.8 所示。Fig. 5.8 的每一行展示一个线程的执行活动，时间从左到右推进。

图中只展示了 block0,0 中线程的活动；其他 block 的行为相同。用于保存 M 元素的 shared memory 数组称为 Mds，用于保存 N 元素的 shared memory 数组称为 Nds。

在 Phase 0 开始时，block0,0 的 4 个线程协作地把一个 M tile 加载到 shared memory 中：

thread0,0 把 M0,0 加载到 Mds0,0
thread0,1 把 M0,1 加载到 Mds0,1
thread1,0 把 M1,0 加载到 Mds1,0
thread1,1 把 M1,1 加载到 Mds1,1

这些加载操作显示在 Fig. 5.8 的第二列。N 的一个 tile 也以类似方式被加载，如 Fig. 5.8 的第三列所示。

这种方式的关键是：每个 phase 只关注输入矩阵中的一个小区域，因此这个小区域可以被放入 shared memory，并被 block 内多个线程重复使用。

4. Shared memory 复用与 locality

在 tiled matrix multiplication 中，shared memory 数组会在不同 phase 中反复使用。

例如书中把保存 M tile 的 shared memory 数组称为 Mds，保存 N tile 的 shared memory 数组称为 Nds。每个 phase 中，Mds 和 Nds 保存当前阶段要使用的输入数据；进入下一个 phase 后，它们会被新的 tile 数据覆盖。

这种做法利用了一个非常重要的概念：locality，局部性。

这里的局部性指的是：程序在某一段时间内，会集中访问一小部分数据。如果这部分数据能够放进更快的内存中，那么大部分访问都可以由高速内存满足，而不需要频繁访问 global memory。

简单说：

没有 tiling：
线程反复从 global memory 读取相同数据

使用 tiling：
线程先把一小块数据加载到 shared memory
然后在当前 phase 中反复使用 shared memory 中的数据

因此，tiling 的本质不是单纯“分块”，而是通过分块制造数据复用机会，把重复访问从 global memory 转移到 shared memory。

5. 本节核心总结

5.3 节主要说明：tiling 是减少 global memory traffic 的重要方法，尤其适合矩阵乘法这种具有大量数据复用的计算。

朴素矩阵乘法的问题是，同一个 M 或 N 元素会被多个线程重复从 global memory 中读取，导致 global memory 带宽成为瓶颈。

tiling 的优化思路是：

把输入矩阵划分成 tile
让 block 内线程协作加载 tile 到 shared memory
在 shared memory 中复用这些数据
用多个 phase 完成完整点积

这样可以显著减少 global memory 访问次数，提高 compute-to-global-memory-access ratio。

这一节也为 5.4 节做铺垫：5.3 主要讲 tiled matrix multiplication 的思想，5.4 会给出具体 CUDA kernel 实现。

5.4 A tiled matrix multiplication kernel

1. 本节目标：把 tiling 思想落实到 CUDA kernel

5.3 节讲的是 tiled matrix multiplication 的思想：让 block 内线程协作，把 M 和 N 的一小块数据从 global memory 加载到 shared memory，然后在 shared memory 中复用这些数据。

5.4 节开始给出具体 kernel 实现。核心代码结构是：

#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;
    int by = blockIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    int Row = by * TILE_WIDTH + ty;
    int Col = bx * TILE_WIDTH + tx;

    float Pvalue = 0;

    for (int ph = 0; ph < Width / TILE_WIDTH; ++ph) {
        Mds[ty][tx] = M[Row * Width + ph * TILE_WIDTH + tx];
        Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * Width + Col];

        __syncthreads();

        for (int k = 0; k < TILE_WIDTH; ++k) {
            Pvalue += Mds[ty][k] * Nds[k][tx];
        }

        __syncthreads();
    }

    P[Row * Width + Col] = Pvalue;
}

这段代码的重点不是语法本身，而是理解：每个 block 负责计算 P 的一个 tile，每个 thread 负责计算其中一个输出元素。

2. shared memory 数组和线程索引的作用

代码中首先声明两个 shared memory 数组：

__shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
__shared__ float Nds[TILE_WIDTH][TILE_WIDTH];

它们分别用来保存当前 phase 中需要使用的 M tile 和 N tile。

因为 shared memory 的作用范围是 一个 thread block，所以每个 block 都会有自己的一份 Mds 和 Nds。同一个 block 内的所有线程都能访问这两个数组。这一点很关键，因为一个线程加载到 shared memory 中的数据，后面可能会被同一个 block 内的其他线程使用。

接着代码把索引保存到局部变量：

int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;

这些变量是普通自动标量变量，通常放在寄存器中。每个线程都有自己的 bx、by、tx、ty。这样写主要是为了代码更简洁，也避免反复写较长的内置变量名。

然后计算当前线程负责的输出元素位置：

int Row = by * TILE_WIDTH + ty;
int Col = bx * TILE_WIDTH + tx;

含义是：

blockIdx.y 决定当前 block 负责 P 的哪一块行区域
blockIdx.x 决定当前 block 负责 P 的哪一块列区域
threadIdx.y 决定 block 内的具体行偏移
threadIdx.x 决定 block 内的具体列偏移

所以每个线程最终负责计算：

P[Row][Col]

在线性内存中就是：

P[Row * Width + Col]

3. phase 循环：分阶段加载输入 tile

核心外层循环是：

for (int ph = 0; ph < Width / TILE_WIDTH; ++ph)

这里的 ph 表示当前是第几个 phase。矩阵乘法中，计算一个 P[Row][Col] 需要遍历完整的 k 维度。tiling 后，完整的 k 维度被切成多个长度为 TILE_WIDTH 的小段，每一段就是一个 phase。

在每个 phase 中，每个线程负责加载一个 M 元素和一个 N 元素到 shared memory：

Mds[ty][tx] = M[Row * Width + ph * TILE_WIDTH + tx];
Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * Width + Col];

对 M 来说，当前线程加载的是：

M 的 Row 行
第 ph 个 tile 中由 tx 决定的列

对 N 来说，当前线程加载的是：

N 的 Col 列
第 ph 个 tile 中由 ty 决定的行

也就是说，整个 block 的线程一起完成：

加载一个 TILE_WIDTH × TILE_WIDTH 的 M tile
加载一个 TILE_WIDTH × TILE_WIDTH 的 N tile

这样，原本每个线程都要单独从 global memory 读的数据，现在由 block 内线程协作加载一次，然后放到 shared memory 中供大家复用。

4. 两个 `__syncthreads()` 的作用

加载完 Mds 和 Nds 后，代码立刻调用：

__syncthreads();

这个同步非常重要。因为 shared memory 中的数据是由整个 block 的线程共同加载的，一个线程后面计算时可能会用到其他线程加载的数据。如果没有同步，有些线程可能还没加载完，另一些线程就已经开始读取 shared memory，结果就可能读到错误数据。

第一个 __syncthreads() 解决的是：

read-after-write dependence

也就是读后依赖：必须等所有线程把数据写入 shared memory 后，其他线程才能读取。

然后每个线程在 shared memory 中做当前 phase 的部分点积：

for (int k = 0; k < TILE_WIDTH; ++k) {
    Pvalue += Mds[ty][k] * Nds[k][tx];
}

这里访问的是 shared memory，不再是 global memory。Pvalue 是每个线程自己的寄存器变量，用来累加最终结果。

完成当前 phase 的计算后，代码再次调用：

__syncthreads();

第二个同步也很重要。因为下一轮 phase 会复用同一块 Mds 和 Nds。如果有些线程提前进入下一轮并覆盖 shared memory，而其他线程还没用完上一轮的数据，就会出错。

第二个 __syncthreads() 解决的是：

write-after-read dependence

也就是写后依赖：必须等所有线程都读完当前 shared memory 中的数据后，才能覆盖它们。

5. 性能收益与本节总结

tiled matrix multiplication 的主要收益是减少 global memory 访问。对于矩阵乘法，使用 tiling 后，global memory accesses 可以减少大约 TILE_WIDTH 倍。

书中举例说，如果使用 32 × 32 tile，global memory accesses 可以减少 32 倍，compute-to-global-memory-access ratio 可以从朴素实现的：

0.25 FLOP/B

提高到：

8 FLOP/B

在 H100 上，若按全局内存带宽 3.35 TB/s 估算，则可支持的计算吞吐量从约：

3.35 TB/s × 0.25 FLOP/B = 0.84 TFLOPS

提高到：

3.35 TB/s × 8 FLOP/B = 26.8 TFLOPS

这说明 tiling 能显著提升性能，但此时 8 FLOP/B 仍低于 H100 的约 20 FLOP/B 阈值，所以这个 kernel 仍可能处在 memory-bound 区域。后续如果想继续逼近峰值性能，还需要更高级的优化，比如更充分的数据复用、更高效的寄存器 tiling、使用 cuBLAS 或 CUTLASS 等优化库。

本节核心可以概括为：

一个 block 计算 P 的一个 tile；
每个线程计算 P tile 中的一个元素；
每个 phase 协作加载一块 M tile 和一块 N tile 到 shared memory；
用 __syncthreads() 保证加载完成和使用完成；
通过 shared memory 复用数据，减少 global memory traffic。

这也是 CUDA 优化中非常经典的模式：用 shared memory 显式管理数据局部性，把重复访存从 global memory 转移到片上存储中。

5.5 Boundary checks

1. 为什么需要 boundary checks

5.4 节中的 tiled matrix multiplication kernel 有一个前提：矩阵宽度 Width 是 TILE_WIDTH 的整数倍，并且矩阵是方阵。在这种情况下，每个 block、每个 thread 访问的 M、N、P 元素都刚好存在。

但实际情况中，矩阵大小不一定刚好能被 tile 大小整除。比如矩阵宽度是 3，而 tile 宽度是 2，那么最后一个 tile 就会越过矩阵边界。此时有些线程会尝试访问不存在的元素。5.5 节就是在解决这个问题：如何让 tiled matrix multiplication kernel 能处理任意尺寸的矩阵。

2. 越界访问会带来什么问题

用 3 × 3 矩阵、2 × 2 tile 举例说明。某些线程在加载 M 或 N tile 时，会访问不存在的元素。

越界访问主要有两类问题。

第一类是访问超过一行末尾的元素。例如某个线程想访问 M0,3，但矩阵一行只有 M0,0、M0,1、M0,2。由于二维矩阵在线性内存中是连续存储的，访问 M0,3 实际可能会读到下一行的 M1,0。这不会一定报错，但会读到错误数据，导致计算结果错误。

第二类是访问整个数组范围之外的地址。例如访问超过矩阵最后一行或最后一列对应的内存位置。这种情况下，程序可能读到随机数据，也可能直接触发错误并终止。

所以问题不是“最后几个线程没用可以不管”，而是：任何一次 global memory 访问，只要索引可能越界，就必须加检查。

3. 不能简单跳过无效输出线程

一个容易想到的做法是：如果某个线程不负责计算有效的 P 元素，那就让它不工作。比如它对应的 Row 或 Col 已经超过矩阵边界，就直接跳过。

但这不够。

原因是：即使某个线程最终不负责写入有效的 P 元素，它仍然可能需要为其他线程加载 M 或 N 的元素到 shared memory。

也就是说，一个线程有三种不同的行为：

加载 M tile
加载 N tile
计算并写回 P

这三种行为的边界条件并不完全一样。因此不能只用一个条件判断“这个线程是否有效”。书中给出的原则是：

每一次 memory access 都要有对应的边界检查，
保证访问使用的索引在被访问数组的合法范围内。

这是本节最关键的工程原则。

4. 加载 M、加载 N、写回 P 的检查方式

在加载 M tile 时，线程要检查它准备访问的 M 元素是否合法。

原来的加载代码类似：

Mds[ty][tx] = M[Row * Width + ph * TILE_WIDTH + tx];

这里访问的二维索引可以理解为：

M 的行索引：Row
M 的列索引：ph * TILE_WIDTH + tx

所以边界检查是：

if (Row < Width && ph * TILE_WIDTH + tx < Width)
    Mds[ty][tx] = M[Row * Width + ph * TILE_WIDTH + tx];
else
    Mds[ty][tx] = 0.0f;

在加载 N tile 时，原来的加载代码类似：

Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * Width + Col];

这里访问的二维索引是：

N 的行索引：ph * TILE_WIDTH + ty
N 的列索引：Col

所以检查条件是：

if (ph * TILE_WIDTH + ty < Width && Col < Width)
    Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * Width + Col];
else
    Nds[ty][tx] = 0.0f;

如果访问越界，就不要从 global memory 中读取，而是往 shared memory 对应位置写入 0.0f。这样做是安全的，因为矩阵乘法里乘以 0 不会改变点积结果。

最后，写回 P 时也要检查当前线程负责的输出元素是否存在：

if (Row < Width && Col < Width)
    P[Row * Width + Col] = Pvalue;

这样可以避免把结果写到非法位置。

5. 本节总结：边界检查让 tiled kernel 更通用

5.5 节的核心是：为了让 tiled matrix multiplication 支持任意矩阵大小，需要对每次 global memory 访问加 boundary checks。

重点可以概括为：

矩阵尺寸不一定是 TILE_WIDTH 的整数倍；
最后一部分 tile 可能越过矩阵边界；
越界读取可能读到错误数据，越界写入可能破坏内存；
不能简单跳过无效输出线程，因为它可能还要为其他线程加载数据；
加载 M、加载 N、写回 P 要分别做边界检查；
越界加载时向 shared memory 写 0.0f。

这一节实际上是在告诉我们：优化代码不能只考虑理想尺寸，还要处理边界情况。5.4 的 tiled kernel 展示了性能优化的基本结构，而 5.5 则让这个结构变得更接近实际可用的通用 kernel。

5.6 Impact of memory usage on occupancy

1. 片上内存能提速，也可能降低 occupancy

5.6 节讨论的是 memory usage 对 occupancy 的影响。前面 5.2 到 5.5 一直在强调 registers 和 shared memory 的好处：它们可以减少 global memory 访问，提高 compute-to-global-memory-access ratio。

但这里要补充一个限制：registers 和 shared memory 都是 SM 上的有限资源。如果一个 kernel 每个线程或每个 block 占用太多片上资源，那么一个 SM 上能同时驻留的线程数量就会下降，occupancy 也会下降。

所以 CUDA 优化里存在一个平衡：

使用更多 registers / shared memory：
    好处：减少 global memory 访问，提高数据复用
    坏处：可能减少 SM 上可驻留线程数，降低 occupancy

使用更少片上资源：
    好处：可能提高 occupancy，增强 latency tolerance
    坏处：可能增加 global memory 访问，降低算术强度

也就是说，shared memory 和 registers 不是越多越好，而是要看它们带来的访存减少，是否值得牺牲 occupancy。

2. Shared memory 如何限制 occupancy

occupancy 受很多因素影响，包括线程数、block 数、register 使用量、shared memory 使用量等。本节重点讲 shared memory。

书中以 H100 为例：H100 每个 SM 最多可以配置 228 KB shared memory，最多支持 2048 个线程/SM。如果想让 2048 个线程都驻留在一个 SM 上，那么平均每个线程可使用的 shared memory 不能超过：

228 KB / 2048 threads = 114 B/thread

对于 tiled matrix multiplication，假设一个 block 有：

TILE_WIDTH × TILE_WIDTH 个线程

并且使用两个 shared memory tile：

__shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
__shared__ float Nds[TILE_WIDTH][TILE_WIDTH];

每个 tile 中有 TILE_WIDTH² 个 float，每个 float 是 4 B。所以每个 block 使用的 shared memory 是：

Mds：TILE_WIDTH² × 4 B
Nds：TILE_WIDTH² × 4 B
总计：2 × TILE_WIDTH² × 4 B

而线程数是：

TILE_WIDTH²

所以平均每个线程使用 shared memory：

(2 × TILE_WIDTH² × 4 B) / TILE_WIDTH² = 8 B/thread

这个数远小于 H100 上满 occupancy 所需的 114 B/thread，因此这个 tiled matrix multiplication kernel 的 occupancy 通常不会被 shared memory 限制。

3. shared memory 使用过多时的例子

书中接着给了一个反例：如果一个 kernel 的每个 block 使用 38 KB shared memory，每个 block 有 256 个线程，那么平均每个线程使用：

38 KB / 256 threads = 152 B/thread

这个值已经超过前面算出的 114 B/thread。于是这个 kernel 无法达到 full occupancy。

H100 每个 SM 有 228 KB shared memory，如果按 152 B/thread 计算，一个 SM 最多只能容纳：

228 KB / 152 B/thread = 1536 threads

而 H100 每个 SM 最多支持 2048 threads，所以最大 occupancy 是：

1536 / 2048 = 75%

这个例子说明：即使线程块大小本身没有问题，只要每个 block 使用 shared memory 太多，也会限制一个 SM 上能同时驻留的 block/thread 数量，从而降低 occupancy。

4. 静态 shared memory 和动态 shared memory

前面 5.4 的 tiled matrix multiplication 使用的是静态 shared memory 声明：

__shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
__shared__ float Nds[TILE_WIDTH][TILE_WIDTH];

这种写法的问题是：shared memory 大小由编译期常量 TILE_WIDTH 决定。如果想改变 tile 大小，就要修改 TILE_WIDTH 并重新编译代码。kernel 无法在运行时根据不同 GPU 的 shared memory 容量灵活调整。

但不同 GPU 的 shared memory 容量可能不同，不同 kernel 场景需要的 tile 大小也可能不同。因此，CUDA 支持 dynamic shared memory。

动态 shared memory 的声明方式类似：

extern __shared__ float Mds_Nds[];

这里不再在 kernel 里写死数组大小，而是在 kernel launch 时通过第三个配置参数指定每个 block 使用多少 shared memory：

size_t size = ...;
matrixMulKernel<<<dimGrid, dimBlock, size>>>(...);

其中 size 的单位是字节。

因为这里只声明了一个一维数组 Mds_Nds[]，所以如果要同时保存 Mds 和 Nds，就需要手动把这块 shared memory 分成两段：

前半段：作为 Mds
后半段：作为 Nds

访问时也不能再直接用二维数组语法，而要用线性化索引手动计算位置。

5. 本节总结：优化要在算术强度和 occupancy 之间取平衡

5.6 节的重点是：片上内存资源可以减少 global memory 访问，但片上资源本身有限，使用过多会降低 occupancy。

可以概括为：

registers 和 shared memory 能提高数据复用，减少 global memory traffic；
但它们都是 SM 上的有限资源；
每个线程或每个 block 使用资源越多，SM 能同时驻留的线程通常越少；
occupancy 降低后，GPU 隐藏访存延迟的能力可能下降；
因此优化时要平衡 arithmetic intensity 和 occupancy。

对 tiled matrix multiplication 来说，当前版本 shared memory 使用量大约是：

8 B/thread

通常不会成为 occupancy 的限制因素。

但对于 shared memory 使用较大的 kernel，例如每个 block 使用几十 KB shared memory，就可能明显降低 occupancy。此时需要结合硬件资源、block size、shared memory 用量和 register 用量一起分析。

本节还引出了一个工程技巧：如果希望 kernel 根据不同硬件或不同问题规模灵活调整 shared memory 用量，可以使用：

extern __shared__

配合 kernel launch 的第三个参数，在运行时动态指定 shared memory 大小。

下面是 5.7 Summary 的中文翻译。

5.7 总结

总的来说，在现代处理器中，一个程序的执行速度取决于 kernel 代码中的 compute-to-global-memory-access ratio，也就是计算与全局内存访问比。

如果这个比值较高，那么 kernel 就是 compute-bound，它的执行速度可以接近硬件的峰值计算吞吐量。

如果这个比值较低，那么 kernel 就是 memory-bound，它的执行速度会受到操作数从内存中被访问出来的速率限制。

CUDA 提供了对 registers、shared memory 和 constant memory 的访问。这些内存比 global memory 小得多，但访问速度要快得多。

把数据放在这些内存中，可以让程序在访问这些数据时不消耗 global memory bandwidth，从而提高代码的 compute-to-global-memory-access ratio。

然而，要有效使用这些内存，需要重新设计算法。书中使用矩阵乘法作为例子，说明了 tiling 这种常见策略。tiling 可以增强数据访问的 locality，也就是局部性，并使 shared memory 能够得到有效使用。

在并行编程中，tiling 会使用 barrier synchronization，也就是屏障同步，强制多个线程在每个执行阶段共同关注输入数据的一个子集。这样，这个数据子集就可以被放入这些特殊内存类型中，从而获得更高的访问速度。

不过，CUDA 程序员必须注意这些特殊内存类型的容量是有限的。它们的容量依赖于具体硬件实现。

一旦使用量超过这些内存的容量限制，就会限制每个 SM 中能够同时执行的线程数量，并可能对 GPU 的计算吞吐量以及隐藏延迟的能力产生负面影响。

因此，在开发应用程序时，能够根据硬件限制进行推理，是并行编程中的一个关键方面。

虽然本章是在 GPU 的 CUDA 编程背景下介绍 tiled algorithm，但它实际上是在几乎所有类型的并行计算系统中实现高性能的有效策略。

原因是：应用程序必须在数据访问中表现出 locality，才能有效使用这些系统中的高速内存。

例如，在多核 CPU 系统中，数据局部性可以让应用程序有效使用片上的 data cache，从而降低内存访问延迟并获得高性能。这些片上 data cache 的容量同样有限，因此也要求计算过程具有局部性。

因此，当读者使用其他编程模型为其他类型的并行计算系统开发并行应用程序时，也会发现 tiled algorithm 很有用。

本章的目标是介绍 locality、tiling 和不同 CUDA memory types 的概念。

本章介绍了一个使用 shared memory 的 tiled matrix multiplication kernel。随后进一步研究了在应用 tiling 技术时，为了支持任意数据维度，需要使用 boundary test conditions，也就是边界测试条件。

本章还简要讨论了 dynamically sized shared memory allocation，即动态大小的 shared memory 分配。这样，kernel 就可以根据硬件能力，调整每个 block 使用的 shared memory 大小。

PMPP 第五章

Memory architechure and data locality

5.1 Memory bandwidth as a performance limiter

1. GPU 性能的两个常见约束

2. Compute-bound、Memory-bound 与算术强度

3. Roofline Model 和 speed-of-light analysis

4. 向量加法：典型 memory-bound kernel

5. 矩阵乘法：理论上 compute-bound，朴素实现却 memory-bound

5.2 CUDA memory types

1. CUDA 内存类型的整体关系

2. Registers：线程私有的最快存储

3. Shared memory：block 内线程共享的高速存储

4. 变量声明方式与内存位置

5. 本节核心总结

5.3 Tiling for reduced memory traffic

1. Tiling 的基本思想：用小而快的内存缓存局部数据

2. 矩阵乘法中的重复访存问题

3. Tiled matrix multiplication 的执行方式

4. Shared memory 复用与 locality

5. 本节核心总结

5.4 A tiled matrix multiplication kernel

1. 本节目标：把 tiling 思想落实到 CUDA kernel

2. shared memory 数组和线程索引的作用

3. phase 循环：分阶段加载输入 tile

4. 两个 __syncthreads() 的作用

5. 性能收益与本节总结

5.5 Boundary checks

1. 为什么需要 boundary checks

2. 越界访问会带来什么问题

3. 不能简单跳过无效输出线程

4. 加载 M、加载 N、写回 P 的检查方式

5. 本节总结：边界检查让 tiled kernel 更通用

5.6 Impact of memory usage on occupancy

1. 片上内存能提速，也可能降低 occupancy

2. Shared memory 如何限制 occupancy

3. shared memory 使用过多时的例子

4. 静态 shared memory 和动态 shared memory

5. 本节总结：优化要在算术强度和 occupancy 之间取平衡

5.7 总结

4. 两个 `__syncthreads()` 的作用