GPU系统拆解-06-内存层次与访存主线：GPU 性能为什么常常卡在数据

本文是「GPU系统拆解」系列第 06 篇。
系列导读：GPU系统拆解-00-导读：从架构认知到推理系统的学习路线
上一篇：GPU系统拆解-05-CUDA 执行映射：抽象模型如何落到真实硬件
下一篇：GPU系统拆解-07-PTX、SASS 与编译链：一段 CUDA 代码如何变成指令

这一篇是整套 GPU 学习里最关键的锚点之一。真正吃透它之后，你再看 GEMM、FlashAttention、PagedAttention、LLM decode、kernel fusion，都会更容易意识到：很多优化的本质不是“把公式算快一点”，而是“把数据放对位置、以对的方式访问，并尽量减少远距离搬运”。

1. 先给结论

• 很多 GPU 性能问题，本质上不是“算不动”，而是“喂不动”。
• GPU 优化的核心，不是机械地多用 shared memory，而是让数据停留在更近、更快、更可复用的层次。
• register、shared memory、cache、global memory 各自解决的是不同层次的问题，不能混着理解。
• coalescing 决定全局访存效率，bank conflict 决定 shared memory 是否真的快，register spill 往往意味着性能风险。
• arithmetic intensity 可以帮助你判断一个算子更像 compute-bound 还是 memory-bound。
• 对 AI infra / 推理 来说，很多难点最终都会落到 memory system：权重读取、KV cache、layout、batching 和数据流组织。

2. 为什么 GPU 特别在意内存层次

CPU 当然也有 cache hierarchy，也有访存瓶颈，但 GPU 对这件事更敏感，原因主要有三个。

2.1 GPU 的吞吐高，数据跟不上就会整体空转

GPU 有大量执行单元。如果数据供应速度不够，空转的不是少数核心，而可能是一大批执行资源。

可以粗略地理解成：

• CPU 更像少量高性能工人
• GPU 更像大规模流水线

流水线最怕的不是“不会干活”，而是“材料没送到”。

2.2 深度学习和推理里有大量 memory-sensitive 场景

很多人会把 AI 理解成纯算力问题，这不完整。

尤其在大模型推理里，经常会出现这种情况：

• 理论 FLOPS 很高
• 但实际 token/s 提升并不理想

根本原因之一是很多阶段并不是纯 compute-bound，而是明显受：

• 权重读取
• KV cache 访问
• cache 命中率
• 访问模式规整性
• 显存带宽

这些因素约束。

2.3 latency hiding 很重要，但不是无限有效

GPU 常靠多 warp 并发来隐藏内存延迟：

• 某个 warp 在等内存
• 调度器切去执行其他 warp

这就是 latency hiding。

但如果：

• 大多数 warp 都在等 memory
• 访存模式很差
• 活跃 warp 数不够
• memory system 已经接近饱和

那么延迟就会暴露出来。

所以 memory hierarchy、occupancy 和 warp 调度，本质上是联动的。

3. 先建立一张 memory hierarchy 脑图

从近到远，可以先建立这样一张脑图：

1. register
2. shared memory
3. L1 / texture / read-only cache
4. L2 cache
5. global memory
6. host memory（如果发生 host-device 交互）

再补两个容易误解的概念：

• local memory
• constant memory

先不要一开始就纠结所有架构细节，先抓住大方向：

• 越靠近计算单元，越快、越小
• 越远，越大、越慢
• 真正的优化关键不是“总往 shared 搬”，而是判断哪些数据该待在哪一层，以及为什么

4. register：最快，但最贵

4.1 register 是什么

register 是线程私有、最快的存储位置。很多局部变量、循环变量和中间结果都会优先放在这里。

比如：

__global__ void add_kernel(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float x = a[idx];
        float y = b[idx];
        c[idx] = x + y;
    }
}

这里的 idx、x、y，通常会尽量被寄存器化。

4.2 为什么它重要

因为它既快，又不需要像 shared memory 那样做线程协作和同步。

很多高性能 kernel 的核心思路之一就是：

• 让线程把自己高频使用的小块数据留在寄存器里
• 减少反复从 shared 或 global 取数

4.3 为什么它也会变成问题

寄存器不是无限的。

如果一个线程使用太多寄存器，那么：

• 一个 SM 能同时驻留的线程 / warp / block 数会下降
• occupancy 可能下降
• 甚至发生 register spill

一旦 spill，本该在寄存器里的值可能落到 local memory 路径，性能通常会明显恶化。

所以这里出现一个非常典型的 GPU 权衡：

更多寄存器通常意味着更高的单线程效率，但也可能意味着更低的并发和更高的 spill 风险。

5. shared memory：最重要的“软件可控快内存”

5.1 shared memory 的本质

shared memory 是 block 内线程共享的高速片上存储。

它最重要的价值不是“比 global 快”，而是：

你可以主动决定把哪些数据搬进来，并让同一个 block 内多个线程反复复用它。

所以它更像一种“软件管理的复用机制”。

5.2 为什么它在 GPU 优化里这么核心

很多高性能 kernel 都在重复同一个套路：

1. 从 global 读一个 tile
2. 搬到 shared memory
3. block 内多个线程重复消费这块 tile
4. 再搬下一块

这就是 tiling 的关键意义。

如果没有 shared memory，同一块数据可能会被多个线程反复从 global memory 读取；有了 shared memory，一次远距离读取就能变成多次近距离复用。

5.3 最经典的例子：矩阵乘

矩阵乘之所以是 GPU 学习主线，不只是因为它重要，还因为它很好地展示了“为什么数据复用能改变性能”。

简化版示意：

__global__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];

    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < N; t += 16) {
        As[threadIdx.y][threadIdx.x] = A[row * N + (t + threadIdx.x)];
        Bs[threadIdx.y][threadIdx.x] = B[(t + threadIdx.y) * N + col];
        __syncthreads();

        for (int k = 0; k < 16; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }

    C[row * N + col] = sum;
}

这段代码最重要的不是语法，而是思路：

• 把数据从 global 搬近
• 让一块数据被更多次使用
• 用更少的远距离访问换更多计算

6. global memory：容量大，但远且贵

6.1 它是什么

global memory 就是 device 主显存空间：

• 数据中心卡上常见是 HBM
• 消费级卡上常见是 GDDR

你通过 cudaMalloc 分配的大部分 device 数据，通常都在这一层。

6.2 它的问题是什么

它的问题不是不能用，而是：

• 延迟远高于片上存储
• 带宽再高也不是无限的
• 访问模式不好时事务效率会明显下降

所以一个 kernel 的性能经常取决于：

• 读写了多少 global memory
• 是否连续访问
• 是否有重复读取
• 是否有机会搬到更近层次做复用

7. local memory 和 constant memory：两个容易误解的概念

7.1 local memory 并不“本地且快”

很多人第一次看到 local memory，会误以为它离线程很近。

这通常是错的。

在 CUDA 语义里，local memory 更接近“线程私有地址空间中的一部分抽象”，它常常出现在：

• register 不够
• 局部数组太大
• 编译器无法把访问很好地寄存器化

时。性能上，它经常沿更慢的 device memory 路径表现。

所以在 profile 里看到大量 local memory traffic，往往不是好消息。

7.2 constant memory 适合广播式读取

constant memory 更适合：

• 数据量不大
• 只读
• 多个线程经常访问同一个位置

它不是通用优化武器，但对一些小型固定参数表很有价值。

8. cache 的作用：L1 / L2 在解决什么问题

初学 CUDA 时，很多人会把注意力全部放在 shared memory 上，容易忽略 cache。

正确理解应该是：

• shared memory 更偏软件主动管理复用
• cache 更偏硬件自动辅助复用

尤其是 L2，对很多工作负载都很关键，因为它会影响：

• 全局流量缓冲
• 热点数据复用
• 某些重复读取场景的有效代价

但不要误以为有 cache 就可以不管访问模式。GPU 的 cache 不会自动替你修复所有糟糕的数据布局和离散访问。

9. coalescing：最关键的全局访存概念

如果只记一个访存关键词，优先记 coalescing。

9.1 它到底是什么

简单说：

同一个 warp 内线程如果访问的是相邻或规整排列的地址，硬件更容易把这些访问合并成较少的内存事务。

这样效率更高。

反过来，如果同一个 warp 里的线程地址很分散，那么：

• 事务更多
• 带宽利用更差
• 访存开销更大

9.2 一个经典对比

连续访问：

__global__ void good_access(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = in[idx];
    }
}

跨步访问：

__global__ void bad_access(const float* in, float* out, int n, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int pos = idx * stride;
    if (pos < n) {
        out[idx] = in[pos];
    }
}

两者功能可以相近，但性能可能差很多，因为同一个 warp 的地址规整性完全不同。

9.3 为什么这对推理特别重要

推理系统里很多 memory-bound 场景都受 coalescing 影响，比如：

• embedding lookup
• KV cache 读取
• gather / scatter
• 分页缓存布局

很多推理框架底层优化，本质上不是“数学更复杂”，而是布局更适合 GPU 访问模式。

10. bank conflict：shared memory 也不是怎么用都快

shared memory 虽然快，但不是无条件快。

它会被划分成多个 bank。如果同一时刻很多线程访问的地址落到同一个 bank，就会产生 bank conflict，导致访问效率下降。

你可以把它理解成：

• 不同 bank 像不同通道
• 访问分散到不同 bank，效率好
• 很多线程挤到同一个 bank，就堵车

为什么矩阵转置常被拿来讲这个问题？因为它经常容易同时遇到：

• global memory 访问规整性问题
• shared memory bank conflict 问题

例如常见的 padding 技巧：

__shared__ float tile[32][33];

核心思想不是“记住 33 这个数字”，而是：

shared memory 的性能也取决于访问模式。

11. arithmetic intensity：把计算和访存统一起来看

这是判断瓶颈类型最实用的概念之一。

11.1 什么是 arithmetic intensity

粗略理解就是：

计算量 / 数据搬运量

如果一个工作负载：

• 做了很多乘加
• 搬的数据相对少

那它通常更接近高 arithmetic intensity，更可能 compute-bound。

反过来，如果：

• 没做多少计算
• 却搬了很多数据

那它通常更接近低 arithmetic intensity，更可能 memory-bound。

11.2 两个最典型的例子

向量加法：

• 每个元素只做一次加法
• 却要读两个数、写一个数
• 算术强度低，通常偏 memory-bound

矩阵乘：

• 一个 tile 搬进来后会被反复做大量乘加
• 数据复用高
• 算术强度更高，更容易逼近 compute-bound

11.3 为什么它对推理很重要

它能帮助你判断优化入口：

• 某个阶段更该从减少访存入手
• 还是更该从提升 Tensor Core 利用率入手

比如：

• prefill 常常更像大矩阵计算，算术强度更高
• decode 常常更像频繁读权重和 KV，memory pressure 更突出

12. 从推理视角看 memory hierarchy：为什么 KV cache 这么麻烦

12.1 KV cache 为什么是 memory system 问题

KV cache 的难点远不只是“存一下历史 token 的中间结果”。

它会直接带来：

• 大容量显存占用
• 高频读取
• 序列变长后的数据膨胀
• 动态请求长度带来的布局不规整

所以从 GPU 角度看，KV cache 真正在问的是：

• 数据放在哪里
• 怎么布局
• 怎么分页
• 怎么减少碎片
• 怎么让读取尽量规整
• 怎么兼顾 batch 动态变化

这也是为什么 PagedAttention 这类设计，本质上首先是内存系统设计问题。

12.2 为什么 decode 更容易 memory-bound

decode 的典型特征是：

• 每步新增计算量有限
• 但要反复读取大量历史状态
• 同时持续读权重和 KV

所以它常常更像：

• 不是“算不完”
• 而是“读不动”

更准确的说法应该是：

decode 阶段通常会暴露更明显的 memory system pressure，因此优化重点往往是权重、KV、cache、layout、batching 的协同，而不只是提高算子 FLOPS。

13. 如何判断一个 kernel 是不是访存问题

这是最实用的分析框架之一。遇到慢 kernel，先问下面六个问题：

13.1 它到底搬了多少数据

• 输入读多少
• 输出写多少
• 是否有中间临时数据
• 是否有重复加载

13.2 数据访问规整吗

• 相邻线程是否访问相邻地址
• stride 是否过大
• 是否存在大量随机 gather / scatter

13.3 数据能复用吗

• 同一个 block 内多个线程会不会用到同一份数据
• 能不能先搬到 shared memory
• 能不能缓存到寄存器

13.4 shared memory 用得合理吗

• 是否有 bank conflict
• 是否 shared memory 占用过大导致 occupancy 降太多
• 是否同步开销太重

13.5 是否发生 spill

• 寄存器是否过多
• local memory traffic 是否异常

13.6 它真的 memory-bound 吗

有时一个 kernel 看起来搬了很多数据，但真正问题也可能是：

• warp divergence 太重
• launch 太碎
• 指令依赖太强
• 计算路径没有被很好向量化或张量化

最终还是要结合 profile，而不是只靠直觉下结论。

14. 常见误区

• 误区 1：GPU 优化主要是把 Tensor Core 用满。
  不够。很多 kernel 首先受 memory system 约束。

• 误区 2：shared memory 一定比 cache 更值得依赖。
  不对。shared memory 是显式复用，cache 是自动复用，价值取决于访问模式。

• 误区 3：local memory 听起来很本地，所以应该很快。
  错。它常常意味着 spill 或较差的存储路径。

• 误区 4：coalescing 只是底层细节，不会决定推理效果。
  错。很多推理热点场景就是被布局和访问模式卡住的。

• 误区 5：高 occupancy 就一定更快。
  错。occupancy 只是潜在并发能力，不是最终性能结果。

15. 本篇必须记住的内容

• 很多 GPU 性能问题，本质上是 memory problem。
• register 最快，但数量最贵；spill 往往是风险信号。
• shared memory 的核心价值是“软件管理的数据复用”。
• global memory 不可避免，但必须尽量少而规整地访问。
• coalescing 决定全局访存效率，bank conflict 决定 shared memory 是否真的快。
• arithmetic intensity 能帮助你判断算子更像 compute-bound 还是 memory-bound。
• 对 LLM 推理来说，KV cache、decode、layout 和 batching 都和 memory hierarchy 深度绑定。

16. 精简版面试表达

shared memory 为什么快

shared memory 快，不只是因为它是片上存储，更重要的是它允许 block 内线程协作复用数据，从而减少对 global memory 的重复访问。所以它真正的价值是“更近 + 可控复用”，而不只是单次访问延迟低。

coalescing 为什么重要

coalescing 的本质是让同一个 warp 内多个线程的全局访存尽量合并成更少的内存事务，从而提高带宽利用率。很多看起来计算不复杂但性能很差的 kernel，本质上就是访存不够 coalesced。

为什么高 occupancy 也可能慢

因为 occupancy 高只说明活跃 warps 多，不代表 memory system 高效。如果访存模式很差、带宽已经饱和、bank conflict 很重，或者存在大量 spill，那么即使 occupancy 高，kernel 也可能仍然很慢。

---

系列导航

• 导读：GPU系统拆解-00-导读：从架构认知到推理系统的学习路线
• 上一篇：GPU系统拆解-05-CUDA 执行映射：抽象模型如何落到真实硬件
• 下一篇：GPU系统拆解-07-PTX、SASS 与编译链：一段 CUDA 代码如何变成指令
• 延伸阅读：CUDA实践-10-访存优化与共享内存冲突
• 延伸阅读：CUDA C编程权威指南：导读