2. GPU内存层次结构：全局内存、共享内存、寄存器、L1/L2缓存的特点与访问延迟
做GPU量化计算优化，说白了就是在跟内存打交道。
我刚开始接触CUDA时，总觉得计算才是核心。后来被现实狠狠教育了一回——一个kernel跑得慢，90%的原因都是内存访问没搞好。你想想看，GPU的计算单元多到吓人，但数据喂不进去，再强的算力也是白搭。
这一章，咱们就把GPU的内存层次结构彻底捋清楚。我会结合自己的踩坑经验，帮你建立直观的“内存延迟感”。
2.1 一张图看懂全局
先放一张我手绘的内存层次结构图。这张图我每次培训都会拿出来讲，因为它能帮你快速建立全局观。
这张图我每次看都有新体会。核心就一句话：离计算单元越近，越快但越小；越远，越慢但越大。量化优化的本质，就是想办法把数据往“上面”挪。
2.2 寄存器：最快但最稀缺
寄存器是GPU里最快的内存，没有之一。延迟只有1个时钟周期，基本就是零延迟。
但问题来了——每个线程能用的寄存器数量非常有限。我记得在Fermi架构时代，每个SM只有32768个寄存器。现在虽然多了不少，但架不住线程数量也爆炸式增长啊。
关键数据：
- 每个线程最多可用255个32位寄存器（不同架构略有差异）
- 一旦超出，编译器会自动做“寄存器溢出”（register spilling）
- 溢出后的数据会被存到全局内存，性能直接崩盘
⚠️ 我曾经踩过的坑：
有一次我写一个量化kernel，用了大量局部变量。编译时没报错，但跑起来慢得离谱。后来用 --ptxas-options=-v 一看，每个线程用了80多个寄存器，发生了严重的溢出。数据被频繁写到全局内存，性能直接掉了10倍。
从那以后，我养成了一个习惯：每次写完kernel，第一件事就是检查寄存器使用量。
2.3 共享内存：手动控制的“神器”
共享内存是我个人最喜欢的GPU内存类型。为什么？因为它可编程、可控制。
L1缓存是硬件自动管理的，你没法干预。但共享内存不一样——你可以精确控制哪些数据放进去，什么时候放，什么时候取走。这种控制力，在量化计算中简直是降维打击。
共享内存的几个关键特点：
- 位于芯片上：延迟约20-30个周期，比全局内存快一个数量级
- 每个SM独立拥有：不同SM的共享内存互不干扰
- 容量有限：通常每个SM只有48KB-164KB（取决于架构）
- 需要手动管理：用
__shared__关键字声明
我举个例子你就明白了。假设你要做矩阵乘法，A和B两个矩阵都在全局内存里。如果直接算，每个元素都要从全局内存读，延迟400-800周期，慢得没法忍。
但如果你把A的一块和B的一块先搬到共享内存里，然后在共享内存里做计算，延迟直接降到20-30周期。这就是量化计算中“分块”（tiling）技术的核心思想。
💡 实用技巧：
共享内存的bank conflict是个大坑。简单说，共享内存被分成32个bank，如果同一个warp里的线程访问同一个bank的不同地址，就会发生冲突，导致串行化。
解决办法？我一般用 __shared__ float s_data[32][32+1] 这种“加一垫”技巧，避免同一列的线程访问同一bank。
2.4 L1/L2缓存：自动但不可控
L1和L2缓存是硬件自动管理的，你没法直接控制。但理解它们的行为，对优化至关重要。
| 特性 | L1缓存 | L2缓存 |
|---|---|---|
| 位置 | 每个SM内部 | 所有SM共享 |
| 容量 | 通常16-128KB（与共享内存共享） | 通常1-6MB |
| 延迟 | ~30 cycles | ~200-400 cycles |
| 管理方式 | 硬件自动 | 硬件自动 |
| 缓存粒度 | 128字节（cache line） | 32字节（sector） |
这里有个重要的点：L1缓存和共享内存共享同一块物理存储。你可以通过 cudaFuncSetAttribute 或编译器选项来分配比例。比如，你可以配成48KB共享内存 + 16KB L1，或者反过来。
我个人习惯：如果kernel对共享内存需求大，就多分给共享内存；如果主要是全局内存的随机访问，就多留点给L1。
2.5 全局内存：容量大但代价高
全局内存就是显存，容量最大（几GB到几十GB），但延迟也最高（400-800 cycles）。
量化计算中，大部分数据一开始都在全局内存里。我们的目标就是：尽量减少对全局内存的直接访问。
全局内存访问有几个关键优化点：
- 合并访问（coalesced access）：同一个warp的线程，尽量访问连续的内存地址。这样硬件可以把多次访问合并成一次大的传输。
- 对齐访问：起始地址最好是128字节的倍数，这样能充分利用cache line。
- 避免随机访问：随机访问会导致大量cache miss，性能直接崩。
一个简单的例子：
// ❌ 非合并访问（慢）
float val = data[threadIdx.x * N + blockIdx.x];
// ✅ 合并访问（快）
float val = data[blockIdx.x * N + threadIdx.x];你看，只是交换了索引顺序，性能就能差出几倍。这就是合并访问的威力。
2.6 延迟对比：一张表看懂
最后，我整理了一张延迟对比表。这张表我建议你打印出来贴在工位上。
| 内存类型 | 延迟（cycles） | 相对比例 | 容量 | 管理方式 |
|---|---|---|---|---|
| 寄存器 | ~1 | 1x | ~256个/线程 | 编译器自动 |
| 共享内存 | ~20-30 | ~20-30x | ~48-164KB/SM | 手动 |
| L1缓存 | ~30 | ~30x | ~16-128KB/SM | 硬件自动 |
| L2缓存 | ~200-400 | ~200-400x | ~1-6MB | 硬件自动 |
| 全局内存 | ~400-800 | ~400-800x | 几GB-几十GB | 手动分配 |
看到这个比例了吗？全局内存比寄存器慢了800倍。这意味着什么？意味着你花1个周期做计算，可能要等800个周期等数据。这就是所谓的“内存墙”。
量化优化的核心，说白了就是打破这堵墙。怎么破？把数据往快了挪——能放寄存器的放寄存器，能放共享内存的放共享内存，尽量减少对全局内存的依赖。
💡 我的个人经验：
做量化计算优化时，我一般遵循这个优先级：
- 能用寄存器解决的，绝不用共享内存
- 能用共享内存解决的，绝不用全局内存
- 实在不行了，再考虑全局内存的合并访问优化
这个原则帮我解决了不少性能问题，你可以试试看。
好了，GPU内存层次结构就讲到这里。记住这张图、这张表，还有我踩过的那些坑。下一章，咱们会深入聊一个具体的优化技术——内存合并访问，到时候你会看到这些理论知识是怎么落地到实际代码里的。