2. 内存层级：合并访问、shared memory 与 bank

这是 GPU 与 CPU 差别最大的一章。CPU 上“写个循环，剩下的交给 cache”，在 GPU 上通常只能拿到峰值带宽的 10%–20%。原因有两个：访问模式要按 warp 对齐，以及shared memory 需要你自己搬。

2.1 五层存储

关键区别不在“快慢”，而在谁负责搬运。CPU 的 cache 对程序透明：你访问一次，硬件自动把整条 cache line 拉进来，下次命中就快。GPU 的 shared memory 是每个 block 独占的 scratchpad：容量固定、生命周期等于 block、必须用 __syncthreads() 协调读写。硬件不会自动把你要用的数据放进去，你要自己写加载代码。

L1 和 shared 在物理上是同一块存储，按 kernel 的需求划分（carveout）。所以“shared 用得越多，L1 越小”，两者不能同时吃满。

2.2 在途字节：为什么 100% 占用率也可能跑不满带宽

回到 0.3 的 Little 定律。要跑满带宽，任意时刻都要有

字节的请求在路上。H100 上约 2 MB；除以 132 个 SM，每个 SM 要维持约 15 KB 的在途数据。

• 一个 warp 用 float 读：一次 128 B → 需要约 119 个 warp 同时等，超过每 SM 64 个的上限。
• 用 float4 读：一次 512 B → 只需要 30 个 warp，在 64 个的限制内。
• 每个线程发两个独立的 float4 请求（ILP）：再减半。

这就是“占用率明明 100%，带宽只有 40%”的最常见原因：每条访存指令携带的字节数太少，或者编译器把独立请求串行化了。先看字节/指令，再看占用率。

2.3 合并访问：一个 warp 的 32 个地址怎么变成事务

显存的读写以 32 B 的 sector 为最小单位，L2 line 是 128 B。一个 warp 的 32 个 lane 同时发出地址，硬件把它们按 sector 归并：

• 连续：32 个 lane 读连续的 32 个 float，覆盖 128 B = 4 个 sector，全部有用 → 1 组事务。这是理想情况。
• 跨步 32：地址相差 128 B，每个 lane 落在不同 sector，搬回 32 × 32 B = 1 KB，只用了 128 B → 8 倍浪费。
• 结构体数组（AoS）：每个元素 16 B，只读其中一个字段，仍然搬 16 B → 4 倍浪费。

coalesce.cu

// 好：一个 warp 覆盖 128 B 连续地址
out[i] = in[i];

// 坏：一个 warp 跨 4 KiB
out[i] = in[i * 32];

// 好：每个 lane 一次搬 16 B（要求 16 B 对齐，且 n 是 4 的倍数）
float4 v = *reinterpret_cast<const float4*>(in + 4 * i);
float4 r;
r.x = v.x * 2.f; r.y = v.y * 2.f; r.z = v.z * 2.f; r.w = v.w * 2.f;
*reinterpret_cast<float4*>(out + 4 * i) = r;

如果数据本来就是“每个元素多个字段”，优先改成结构体数组（SoA）：

// AoS：读 x 时每个 lane 拉 16 B，只用 4 B
struct Particle { float x, y, z, w; };
Particle* p;

// SoA：三个独立数组，读 x 完全合并
float *px, *py, *pz, *pw;

2.4 间接寻址：当地址是查表来的

2.3 的三种形态有一个共同前提：地址由下标算出来，i 定了，地址就定了。分页 KV 打破了这个前提 —— 逻辑上的第 3 页在物理上可能是第 42 页，地址要先读一张表。

逻辑页 0  →  物理页 17   ┐
逻辑页 1  →  物理页 3    │  连续的逻辑页
逻辑页 2  →  物理页 42   │  物理上不连续
逻辑页 3  →  物理页 8    ┘
           每页 128 token：页内仍然连续

图里最重要的一点是：页内仍然是连续的。所以 2.3 的账没有被推翻，只是换了一个变量。

账没有变。最小单位还是 32 B sector，浪费仍然是「每 lane 有用字节 ÷ 32 B」。页表改变的不是公式，而是「32 个 lane 的地址落在几个地方」：

• 每 lane 4 B（一个 float）：32 个 lane 只要 128 B，实际搬回 1024 B → 8 倍。
• 每 lane 16 B（float4）：只要 512 B，搬回 1024 B → 2 倍。
• 每 lane 528 B（一条 token 的 KV，单层）：要 16.5 KB，搬回 17 KB → 1.03 倍，只剩 sector 取整。

看最后一条：每个 lane 一次要的字节够大时，sector 取整的损失只剩 3%。「分页 KV 很浪费带宽」这句话在字节层面基本不成立 —— 真正的成本在下面三条。

1. 依赖加载。block_table[i] 要先读一次，数据地址依赖它的结果，编译器不能提前发这条请求。2.2 的在途字节靠的正是「独立请求能排队」，依赖链把这个 ILP 砍掉，表现就是「占用率够、带宽上不去」。

2. 页边界。一个 warp 读 512 B，起始位置在页内随机，跨页概率约 512 ÷ 页字节。以 KV 为例（每 (层 · token) 528 B）：16 token 一页（8448 B）→ 6.1%；128 token 一页（67584 B）→ 0.76%。差 8 倍，但一次跨页只多一组事务 —— 这是三条里最小的一项，别把它当主因。

3. 并行度。32 个 lane 各落在不同页时，硬件没法把它们合并成一条 128 B line：每个 lane 一个 sector，同时在飞的请求数涨上去，而能用来藏延迟的 warp 只有 64 个（1.5）。页表本身不贵，贵的是它把「一次请求」变成「32 次」。

三个缓解。

• 调页大小。16 → 128 token，跨页概率降 8 倍；代价是尾页（平均半个）：4096 token 的序列，128 页的尾浪费约 1.6%，16 页只有 0.2%。页大小是在「跨页概率」和「尾页浪费」之间取，不是越大越好。
• 把页号提到 warp 维度。让一个 warp 只处理同一页里的 token（page-major 排布），页就成了合并的单位，32 个 lane 又回到一条 line。
• 整页搬进 shared 再挑。顺序读一整页 → shared → 在 shared 里按索引取。随机访问的代价从「显存 sector」换成「shared bank conflict」（2.5），量级差一个数量级。生产框架里的 KV 加载接口就是按「连续的物理页区间」批量搬，正是这个形状。

极端形态：稀疏索引。decode 时常常不是读整个窗口，而是按 top-k 索引读 KV（某生产框架的 decode 直接拿 indices_in_kvcache 做注意力）。这时 32 个 lane 的 32 个 token 可能来自 32 个不同页，而每个 lane 只取其中 4 B 的一个分量 —— 回到 8 倍。同一个公式，粒度不同，答案就不同：先看每 lane 要多少字节，再看地址怎么来。

练习。取 block_size 为 16 和 128 两种页，各画出一个 warp 读 512 B 的三种落点（页内、刚好跨页、跨两页），算出事务数与有效字节比；再回答：为什么 128 页更划算，但生产里不会无限调大？

2.5 shared memory 与 bank conflict

shared memory 被划分成 32 个 bank，每个 bank 宽 4 B。一次访问中，如果 warp 的 32 个地址落在 32 个不同 bank 上，就能一次完成；如果多个 lane 落在同一个 bank 的不同地址，就串行化（n-way conflict）；如果落在同一地址，则是广播，不冲突。

bank 编号就是字地址对 32 取模：

于是 float tile[32][32] 有一个经典陷阱：按列访问时，元素 tile[0][c]、tile[1][c]…的地址相差 32 个字，全部落在同一个 bank → 32 路冲突，慢 32 倍。把第二维加 1：

__shared__ float tile[32][33];   // 每行多一个填充字
// bank(tile[r][c]) = (33r + c) mod 32 = (r + c) mod 32
// 连续变化的 r 或 c 都会落到不同的 bank

填充一个 float 只多花 128 B，却能消掉整个冲突。这是 GPU 上最划算的优化之一。

2.6 实战：矩阵转置

转置是“读合并、写就必然不合并”的典型。朴素写法让一个方向完美、另一个方向完全散开：

transpose.cu

constexpr int T = 32;

// 朴素：读合并（x 连续），写跨 rows 个元素 → 不合并
__global__ void transpose_naive(const float* in, float* out, int rows, int cols) {
    int x = blockIdx.x * T + threadIdx.x;   // 输入列
    int y = blockIdx.y * T + threadIdx.y;   // 输入行
    if (x < cols && y < rows)
        out[x * rows + y] = in[y * cols + x];
}

// 分块：全局读写都合并，转置在 shared 里完成
__global__ void transpose_tiled(const float* in, float* out, int rows, int cols) {
    __shared__ float tile[T][T + 1];        // +1 消除 bank conflict

    int x = blockIdx.x * T + threadIdx.x;
    int y = blockIdx.y * T + threadIdx.y;
    if (x < cols && y < rows)
        tile[threadIdx.y][threadIdx.x] = in[y * cols + x];
    __syncthreads();

    int row = blockIdx.x * T + threadIdx.y; // 输出行 = 输入列
    int col = blockIdx.y * T + threadIdx.x; // 输出列 = 输入行
    if (row < cols && col < rows)
        out[row * rows + col] = tile[threadIdx.x][threadIdx.y];
}

账很容易算：转置的流量是读 4 B + 写 4 B = 8 B/元素。1 GB 数据就是 8 GB 流量，在 3.35 TB/s 下约 2.4 ms。朴素版的写端每次浪费 8 倍 sector，读端正常，所以总时间大约是理想值的 4–5 倍；分块版把浪费消掉，代价只是 4 KB 的 shared 和一次 barrier。

顺便注意：transpose_tiled 的读端仍然是一个 warp 读连续 32 个 float（因为 threadIdx.x 最快变化），写端也是。转置的“不连续”被关在了 shared memory 里，而 shared 的随机访问代价只有 bank conflict，没有 sector 浪费。

2.7 只读数据与向量化

只读的输入加上 const __restrict__，编译器才会放心走只读缓存（LDG），不会因为“可能和输出别名”而插入额外同步：

__global__ void axpy(float* __restrict__ y,
                     const float* __restrict__ x, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

L2 是全卡共享的，跨 block 的复用发生在 L2。矩阵乘里 A、B 被不同 block 反复读，靠的就是 L2 命中；再往上才是 shared 的 block 内复用（第 5 节）。

2.8 和 CPU 的对照

• 缓存管理：CPU 自动；GPU 的 shared 要手写加载/同步，且只活在一个 block 内。
• 最小粒度：CPU cache line 64 B；GPU sector 32 B、L2 line 128 B。跨步访问在两边都亏，但 GPU 的惩罚更直接（按 sector 计费）。
• 隐藏延迟：CPU 靠乱序执行 + 硬件预取；GPU 靠大量 warp 轮换 + ILP，没有预取器。
• 一致性：CPU 多核有缓存一致性协议；GPU 的 L1 不跨 SM 一致，跨 block 通信必须用原子或 fence（3.3）。
• 分配：malloc 便宜、可超卖、可换页；cudaMalloc 同步且慢（驱动要建页表映射），显存不可超卖，OOM 是硬失败。

2.9 容易写错的地方