2、GPU架构与显存层次：NVIDIA GPU架构概览，显存层次结构（HBM、L2 Cache、L1 Cache），显存带宽与延迟
聊到吞吐压测，你绕不开GPU的显存。说白了，压测就是在跟显存打交道。我见过太多人，拿着A100跑模型，结果吞吐上不去，第一反应是“代码写得不好”。其实很多时候，是你不了解GPU的“脾气”。
今天我们就来扒一扒NVIDIA GPU的架构，尤其是显存这块。嗯，这部分内容有点硬核，但相信我，搞懂了它，你调优的时候会少走很多弯路。
NVIDIA GPU架构概览：从Fermi到Hopper
先看一张图，这是我个人习惯的总结方式——把GPU架构拆成几个核心模块：
这张图我画了好几次才满意。你注意看，数据是从HBM一层层往上“喂”到SM的。每一层都有它的脾气——容量、带宽、延迟都不一样。
从Fermi架构开始，NVIDIA就在不断优化这个层次结构。到了Hopper（H100），HBM3的带宽已经飙到3.35TB/s了。但你知道吗？我去年调一个推理服务，发现实际有效带宽连理论值的60%都没跑到。问题出在哪？后面会讲。
显存层次结构：HBM、L2 Cache、L1 Cache
咱们一层层拆开看。
1. HBM（高带宽显存）
HBM是GPU的“大仓库”。它容量大，但离计算单元最远。A100用的是HBM2e，40GB或80GB版本。H100升级到了HBM3。
这里有个关键点：HBM的带宽虽然高，但延迟也高。我记得有一次压测一个BERT模型，发现GPU利用率只有30%。一查，发现大部分时间都在等数据从HBM搬过来。这就是典型的“显存带宽瓶颈”。
核心指标：
- A100 (HBM2e): 带宽 2.0 TB/s，容量 40/80 GB
- H100 (HBM3): 带宽 3.35 TB/s，容量 80 GB
- V100 (HBM2): 带宽 0.9 TB/s，容量 16/32 GB
2. L2 Cache
L2 Cache是所有SM共享的“中转站”。A100的L2 Cache有40MB，H100有50MB。你想想看，如果数据能命中L2，就不用去HBM拿了，延迟能降一个数量级。
我做过一个实验：把矩阵乘法中的分块大小从256调到512，L2命中率从70%升到了92%。吞吐直接涨了1.8倍。这就是L2的威力。
调优技巧： 尽量让数据“热”在L2里。比如做Attention时，把Q、K、V的tile大小控制在L2容量以内。我曾经用这个技巧，把GPT推理的延迟从12ms降到了7ms。
3. L1 Cache / 共享内存
L1 Cache是每个SM私有的，A100每个SM有192KB（可配置为L1或共享内存）。它离CUDA核心最近，延迟只有几十个时钟周期。
但容量太小了。你放不下大矩阵，只能放一些“热数据”。比如做卷积时，把输入feature map的一个小窗口放到共享内存里，反复用。
嗯，这里要注意：L1和共享内存是“共享”同一块物理空间的。你可以通过编译器选项调整它们的比例。我一般习惯把共享内存配大一点，因为很多kernel需要用它做数据交换。
显存带宽与延迟：为什么你的压测跑不满？
好，到了最核心的部分。显存带宽和延迟，是压测的“命门”。
先看一组数据：
| 层级 | 容量 | 带宽 | 延迟（近似） |
|---|---|---|---|
| HBM | 40-80 GB | ~2 TB/s | ~400 ns |
| L2 Cache | 40-50 MB | ~4 TB/s | ~200 ns |
| L1 Cache | 128-192 KB | ~12 TB/s | ~30 ns |
| 寄存器 | 256 KB/SM | ~20 TB/s | ~1 ns |
你看，从HBM到寄存器，带宽差了10倍，延迟差了400倍。这就是为什么我们要做“数据局部性”优化——尽量让数据待在离计算单元近的地方。
我曾经踩过一个坑：跑一个大规模推荐模型，batch size设得太大，导致L2 Cache被频繁“踢出”。结果有效带宽从1.5TB/s掉到了0.6TB/s。后来我把batch size减半，并用了梯度累积，吞吐反而上去了。
避坑指南： 不要盲目追求大batch size。它会让你的数据“溢出”L2 Cache，导致频繁访问HBM。我曾经因为这个原因，一个模型训练时间从2天拖到了5天。后来用nsys profile一看，L2命中率只有40%。
那怎么测有效带宽呢？我一般用NVIDIA的bandwidthTest工具，或者自己写一个简单的kernel：
// 一个简单的带宽测试kernel
__global__ void bandwidth_test(float* in, float* out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
out[idx] = in[idx] * 2.0f; // 简单的读-写操作
}
}
// 调用方式
int N = 1 << 28; // 256M个元素
float *d_in, *d_out;
cudaMalloc(&d_in, N * sizeof(float));
cudaMalloc(&d_out, N * sizeof(float));
// 计时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
bandwidth_test<<<N/256, 256>>>(d_in, d_out, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
float bandwidth = (2.0f * N * sizeof(float)) / (ms / 1000.0f) / 1e9; // GB/s
printf("Effective bandwidth: %.2f GB/s\n", bandwidth);
这个代码很简单，但能反映出真实的有效带宽。我一般会跑三次取中位数，避免冷启动的影响。
最后说一句：显存层次不是死的。你可以通过调整数据布局、访问模式来“欺骗”硬件，让数据尽量留在L1或L2里。这就是压测调优的精髓。
想要更多资料，联系可可