一、显存基础与瓶颈分析
做GPU优化这么多年，我最大的感触就是——显存管理是性能的命门。很多同学一上来就盯着计算量，觉得算得越快越好。但实际项目中，我见过太多计算效率很高、却被显存带宽卡死的案例。说白了，GPU是个"搬运工"，数据搬得快，算得才有意义。
1.1 GPU显存层次结构
先聊聊显存的"家族成员"。GPU里的存储系统，其实是个金字塔结构。我习惯把它分成四层：
- Register（寄存器）：离计算单元最近，延迟几乎为0。但数量极其有限，每个线程通常只有几十到几百个。我在优化一个矩阵乘法时，就因为寄存器溢出，性能直接腰斩。
- Shared Memory（共享内存）：同一个Block内的线程共享，延迟约20-30个周期。容量小，通常几十KB。但用好它，性能能翻倍。
- Local Memory（局部内存）：寄存器不够用时，数据会被"挤"到这里。其实它物理上在Global Memory里，延迟很高。嗯，这是个坑，后面会细说。
- Global Memory（全局内存）：容量最大（几GB到几十GB），但延迟最高（400-800个周期）。带宽虽然大，但用不好就是瓶颈。
核心观点：优化的本质，就是让数据尽量往金字塔上层走。能放寄存器的，别放共享内存；能放共享内存的，别放全局内存。
我画了一张图，帮你直观理解这个层次结构：
1.2 显存带宽与延迟的量化分析
你想想看，为什么有时候计算单元在"空转"？因为数据没到。这就是带宽和延迟的问题。
带宽：单位时间内能搬运多少数据。比如NVIDIA A100，带宽约1.6TB/s。听起来很大对吧？但一个简单的向量加法，如果数据在Global Memory里，带宽利用率可能只有20%。
延迟：从发出请求到拿到数据的时间。Global Memory的延迟是400-800个周期。这段时间里，SM（流多处理器）能执行上千条指令。所以，隐藏延迟是优化的关键。
| 存储层次 | 典型延迟 | 典型带宽 | 容量 | 作用域 |
|---|---|---|---|---|
| Register | ~1 cycle | 极高 | 255个/线程 (Max) | 单个线程 |
| Shared Memory | ~20-30 cycles | ~1-2 TB/s | 48-164 KB/Block | Block内线程 |
| Local Memory | 400-800 cycles | 同Global | 受限于显存 | 单个线程 |
| Global Memory | 400-800 cycles | ~0.9-2 TB/s | GB级别 | 所有线程 |
我的经验：判断一个kernel是否带宽受限，有个简单方法——算一下"计算量/访存量"（即算术强度）。如果这个值小于某个阈值（比如10 FLOPs/Byte），那基本就是带宽瓶颈。我在优化BERT推理时，就发现很多层算术强度只有2-3，果断转向访存优化。
1.3 典型显存瓶颈的Profiling方法
光靠猜可不行。我习惯用工具说话。这里分享几个我常用的方法：
方法一：ncu（NVIDIA Nsight Compute）
这是我最常用的profiler。它能告诉你：
- Memory Throughput：实际带宽利用率
- L1/TEX Cache Throughput：缓存命中情况
- Shared Memory Bank Conflicts：共享内存冲突
- Register Spill：寄存器溢出情况
举个例子，跑一个简单的kernel：
// 一个典型的带宽受限kernel
__global__ void vector_add(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}用ncu跑一下：
ncu --set full -o vector_add ./vector_add看报告里的Memory Throughput。如果远低于理论带宽（比如只有30%），那说明有优化空间。我曾经在一个项目中，通过合并内存访问，把带宽利用率从25%提升到了85%。
方法二：检查Register Spill
这个坑我踩过。当你给每个线程分配太多局部变量时，寄存器不够用，数据会被"挤"到Local Memory。你想想看，本来1个周期能拿到的数据，现在要等400个周期，性能不崩才怪。
怎么查？用--ptxas-options=-v编译选项：
nvcc -o vector_add vector_add.cu --ptxas-options=-v输出里会显示regcount和lmem使用量。如果lmem大于0，就要警惕了。
我曾经踩过的坑：有一次优化一个图像处理kernel，为了减少Global Memory访问，我把所有中间结果都存成局部变量。结果寄存器溢出严重，性能反而下降了30%。后来我改用Shared Memory做缓存，才把性能提上来。记住：不是所有局部变量都适合放寄存器。
方法三：分析内存访问模式
Global Memory的带宽利用率，很大程度上取决于合并访问（Coalesced Access）。简单说，就是让同一个warp里的线程，访问连续的内存地址。
用ncu的Memory Workload Analysis页面，可以看：
- Global Load/Store Efficiency：理想值是100%
- Sector Utilization：越高越好
如果效率低，通常是因为：
- 线程访问地址不连续（比如按列访问矩阵）
- 数据类型对齐有问题
- 存在Bank Conflict（共享内存）
一句话总结：显存优化的核心就三件事——减少访存量、提高带宽利用率、隐藏访存延迟。后面的章节，我会逐一展开这些技巧。
好了，这一章先到这里。记住，没有profile就没有优化。拿到一个新kernel，第一件事就是跑ncu，看看瓶颈在哪。别像我早期那样，闷头优化半天，结果发现瓶颈根本不在我以为的地方。