3、内存层次结构：Cache层级、HBM、GDDR、DDR、共享内存、本地内存
聊异构计算，绕不开内存。说白了，数据在芯片里怎么跑、跑多快，全看内存层次怎么搭。我刚开始接触异构平台时，总觉得计算单元才是主角，后来被一次性能瓶颈狠狠教育了一顿——才发现，内存才是真正的命门。
今天咱们就把这个层次结构掰开揉碎。你想想看，一个数据从硬盘到CPU/GPU核心，中间要经过多少级缓存、多少种内存？每一级都有它的脾气。
3.1 为什么需要这么多层？
原因很简单：快的东西贵，便宜的东西慢。如果全用SRAM做内存，芯片面积会大到离谱，成本直接上天。如果全用DDR，那计算单元大部分时间都在等数据。
所以业界搞了个折中方案——金字塔结构。越靠近计算核心的，越快、越小、越贵；越远的，越慢、越大、越便宜。
核心原则：让最常用的数据待在最快的地方，不常用的数据扔到慢的地方去。这就是局部性原理在硬件上的体现。
3.2 Cache层级：速度的缓冲带
Cache这东西，我习惯把它叫作「计算单元的临时工位」。CPU和GPU都有多级缓存，但设计思路不太一样。
3.2.1 CPU的Cache层级
典型的CPU有L1、L2、L3三级缓存。L1最快，通常分指令缓存和数据缓存，每个核心独享。L2也是每个核心独享，但容量大一些。L3是所有核心共享的。
| 层级 | 典型大小 | 延迟（周期） | 归属 |
|---|---|---|---|
| L1 | 32KB - 64KB | 3-5 | 每核心独享 |
| L2 | 256KB - 1MB | 10-20 | 每核心独享 |
| L3 | 8MB - 64MB | 30-50 | 所有核心共享 |
我在项目中遇到过一个问题：某个多线程程序在CPU上跑，性能始终上不去。后来发现是伪共享（False Sharing）搞的鬼——两个线程各自修改不同变量，但这两个变量恰好落在同一个Cache Line里，导致L1缓存不断失效。解决办法很简单，把变量对齐到64字节边界就行。
避坑指南：我曾经因为没注意Cache Line对齐，白白浪费了两周调优时间。记住，x86的Cache Line一般是64字节，ARM有些是64或128字节。写数据结构时，把热点数据凑在一起，别让无关数据挤进来。
3.2.2 GPU的Cache层级
GPU的缓存结构跟CPU差别很大。NVIDIA的GPU有L1和L2两级缓存。但L1在较新架构（比如Volta之后）里和共享内存共用一块物理空间，可以灵活配置。
举个例子，在A100上，每个SM有192KB的SRAM，你可以配成L1和共享内存的任意比例。我个人习惯把计算密集型的任务配成更多共享内存，把访存密集型的任务配成更多L1。
3.3 HBM：高带宽内存的王者
HBM（High Bandwidth Memory）是GPU和高性能计算领域的明星。它通过硅中介层把多个DRAM die堆叠在一起，再用TSV（硅通孔）连接。带宽高得吓人——HBM2e单颗就能到460GB/s，HBM3更是奔着1TB/s去了。
为什么HBM这么快？因为它位宽大。普通DDR的位宽是64位，HBM的位宽能到1024位。你想想看，同样频率下，位宽翻16倍，带宽就翻16倍。
关键点：HBM的延迟其实比DDR高一些（因为要经过中介层），但带宽优势太明显了。所以它特别适合那些需要大量连续数据吞吐的场景，比如深度学习训练、科学计算。
我记得有一次做大规模矩阵乘法优化，用HBM的GPU比用GDDR的GPU快了将近3倍。原因很简单——矩阵乘法对带宽极度敏感，HBM的带宽优势直接转化成了性能优势。
3.4 GDDR：图形内存的老将
GDDR（Graphics Double Data Rate）是专门为显卡设计的内存。从GDDR5到GDDR6X，每一代都在提升频率和带宽。GDDR6X甚至用了PAM4信号编码，一个时钟周期传4个bit。
GDDR和HBM的区别在哪？说白了，GDDR走的是「高频路线」，HBM走的是「宽位路线」。GDDR频率高、位宽相对小（通常384位或512位），HBM频率低、位宽极大。
| 类型 | 典型带宽 | 位宽 | 典型应用 |
|---|---|---|---|
| GDDR6 | ~500 GB/s | 384位 | 消费级显卡 |
| GDDR6X | ~1 TB/s | 384位 | 高端显卡（RTX 3090/4090） |
| HBM2e | ~460 GB/s per stack | 1024位 | 数据中心GPU、AI加速卡 |
| HBM3 | ~1 TB/s per stack | 1024位 | 最新数据中心GPU |
注意：GDDR的功耗比HBM高。HBM因为堆叠设计，数据传输距离短，功耗更低。所以在功耗受限的场景（比如笔记本、嵌入式设备），HBM更有优势。
3.5 DDR：系统内存的基石
DDR（Double Data Rate）是系统内存的主流。从DDR4到DDR5，每一代都在提升频率和降低电压。DDR5的带宽能到~50 GB/s，虽然跟HBM没法比，但胜在容量大、成本低。
在异构计算中，DDR通常作为CPU的主存，GPU通过PCIe总线访问它。但这里有个坑——PCIe的带宽远低于GPU本地内存。所以数据最好提前搬到GPU的HBM或GDDR里，别在运行时频繁通过PCIe搬运。
我曾经见过一个项目，工程师把数据放在DDR里，每次计算都通过PCIe读取。结果GPU利用率不到10%，全在等数据。后来改成预取到HBM里，利用率直接飙到90%以上。
3.6 共享内存与本地内存：GPU的私房钱
共享内存（Shared Memory）和本地内存（Local Memory）是GPU编程里最需要关注的两个概念。
3.6.1 共享内存
共享内存是GPU每个SM内部的一块SRAM，所有线程块（Thread Block）里的线程都能访问。它的速度跟L1缓存差不多，但由程序员显式控制。
为什么需要共享内存？因为全局内存（Global Memory）太慢了。如果多个线程需要反复访问同一块数据，把它搬到共享内存里，能省下大量时间。
举个例子，做矩阵分块乘法时，我会把子矩阵先搬到共享内存里，然后再计算。这样每个数据只从全局内存读一次，后面都在共享内存里操作。
// CUDA示例：使用共享内存做矩阵分块乘法
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
float sum = 0.0f;
for (int tile = 0; tile < N / BLOCK_SIZE; ++tile) {
// 协作加载数据到共享内存
As[ty][tx] = A[by * BLOCK_SIZE + ty][tile * BLOCK_SIZE + tx];
Bs[ty][tx] = B[tile * BLOCK_SIZE + ty][bx * BLOCK_SIZE + tx];
__syncthreads();
// 从共享内存计算
for (int k = 0; k < BLOCK_SIZE; ++k)
sum += As[ty][k] * Bs[k][tx];
__syncthreads();
}
C[by * BLOCK_SIZE + ty][bx * BLOCK_SIZE + tx] = sum;
}经验之谈：共享内存的大小是有限的。NVIDIA的SM通常有48KB到164KB的共享内存（取决于架构）。别一股脑把所有数据都塞进去，算好每个线程块需要多少，别超了。
3.6.2 本地内存
本地内存听起来像「本地」的，但其实它存在全局内存里，只是每个线程私有的。编译器会把寄存器放不下的局部变量自动 spill 到本地内存。
嗯，这里要注意：本地内存的延迟跟全局内存一样高。如果你发现内核里用了很多局部变量，而且性能不对劲，可以检查一下是不是寄存器溢出导致大量本地内存访问。
我有个习惯：写GPU内核时，尽量控制每个线程的寄存器使用量。NVIDIA的编译器选项 --maxrregcount 可以限制寄存器数量，但设得太低会导致 spill 到本地内存，反而更慢。需要找到那个平衡点。
3.7 一张图看懂内存层次
下面这张SVG图，把整个内存层次结构串起来了。从最靠近计算核心的寄存器，到最远的DDR/SSD，每一级的容量、延迟、带宽都标出来了。
3.8 实战中的选择策略
讲了这么多理论，最后聊聊实际中怎么选。我总结了几条经验：
- 数据复用率高 → 用共享内存或L1缓存。比如卷积、矩阵乘法里的权重数据。
- 数据量大且连续 → 用HBM或GDDR。比如视频帧、大数组。
- 数据量小且随机 → 用寄存器或L1。比如查找表、临时变量。
- 跨设备共享 → 用DDR + PCIe，但尽量批量传输，别小包频繁搬运。
最后提醒一句：别迷信理论带宽。实际能达到的带宽往往只有理论值的60%-80%。原因很多——地址冲突、Bank冲突、访问模式不连续等等。写代码时，尽量让访存模式是连续的、对齐的。
好了，内存层次结构就聊到这儿。下一节咱们会深入Cache一致性协议，看看多核之间怎么保证数据不错乱。那个话题更有意思。