第1章：内存访问模式分析
大家好，我是你们的GPU系统优化讲师。今天咱们聊聊内存访问模式——这玩意儿，说白了就是决定你GPU性能能不能跑满的关键。我见过太多同学，kernel写得花里胡哨，结果一跑profile，内存带宽利用率不到20%。嗯，今天我们就把它彻底搞明白。
1.1 合并访问（Coalesced Access）原理
先问个问题：为什么GPU显存带宽那么大，但你的程序就是跑不快？
答案往往出在「合并访问」上。GPU的显存控制器，一次读取可不是一个数据，而是一整块——通常是32字节、64字节或128字节。如果你让同一个warp里的32个线程，访问连续的内存地址，那GPU就能把这些请求合并成一次或少数几次内存事务。这就是合并访问。
核心原则：同一个warp内的线程，访问地址要连续、对齐。这样显存带宽才能跑满。
举个例子。假设每个线程要读一个float（4字节），32个线程就是128字节。如果地址连续，一次128字节的事务就搞定了。但如果地址是乱序的，比如线程0读地址A，线程1读地址A+100，那GPU就得发起多次小事务，带宽利用率直线下降。
我在项目中遇到过这样一个案例：一个图像处理kernel，原本带宽利用率只有15%。我一看代码，发现每个线程处理的是隔列的像素——说白了就是跨步访问。改成连续行访问后，带宽利用率直接飙到85%。你说这差距大不大？
1.2 对齐与跨步访问
对齐，指的是内存地址要落在某个边界上。比如128字节对齐，就是说地址是128的整数倍。为什么重要？因为GPU的内存控制器是按「段」来读取的。如果地址不对齐，一个请求可能跨越两个段，那就得两次事务，白白浪费带宽。
跨步访问就更常见了。比如你有一个二维数组，按列访问——每个线程读同一列的不同行。这时候，线程0读A[0][col]，线程1读A[1][col]……地址间隔是行宽（stride）。如果行宽不是1，那地址就不连续了。
避坑指南：我曾经在一个矩阵转置的kernel里，按列写入结果。结果profile一看，写带宽利用率只有5%。改成按行写入后，问题迎刃而解。记住：写操作同样需要合并访问！
怎么判断是否对齐？我个人的习惯是，用cudaMemGetInfo查看分配地址，或者直接用__align__属性强制对齐。对于动态分配的内存，可以用cudaMallocPitch来保证行对齐。
1.3 Bank Conflict（Shared Memory）
Shared Memory是GPU上的高速缓存，但用不好反而会拖慢速度。原因就是Bank Conflict。
Shared Memory被分成32个Bank（在较新架构上是32个），每个Bank宽度4字节。如果同一个warp里的多个线程，同时访问同一个Bank的不同地址，就会发生Bank Conflict——这些访问会被串行化，一次只能处理一个。
举个例子：
// 假设 shared_mem[32][32] 是 float 类型
int tid = threadIdx.x;
float val = shared_mem[tid][tid]; // 每个线程访问不同行、不同列
// 这里没有Bank Conflict，因为每个线程访问的是不同Bank但如果是这样：
float val = shared_mem[tid][0]; // 所有线程访问同一列
// 这里就有Bank Conflict了！因为同一列的数据在同一个Bank里为什么会这样？因为Shared Memory的地址映射是：地址 % 32 决定它落在哪个Bank。如果所有线程访问的地址模32都相同，那就全挤在一个Bank里了。
我的经验：解决Bank Conflict最常用的方法是「填充」——在声明shared memory时，多分配一个维度。比如 __shared__ float smem[32][33]; 而不是 [32][32]。多出来的那一列，让地址映射错开，Bank Conflict就消失了。
你想想看，如果每个Bank的带宽是100GB/s，32个Bank就是3.2TB/s。但一旦发生Bank Conflict，实际带宽可能降到100GB/s。这损失，谁受得了？
1.4 Profiling工具解读
光说不练假把式。我们得用工具说话。我个人最常用的是NVIDIA Nsight Compute和nvprof（老项目里还有用）。
关键指标看这几个：
| 指标 | 含义 | 理想值 |
|---|---|---|
| Memory Throughput | 显存带宽利用率 | 接近理论峰值 |
| L1/Shared Memory Throughput | L1或Shared Memory的带宽 | 越高越好 |
| Global Load/Store Efficiency | 合并访问效率 | 100% |
| Bank Conflicts | Shared Memory Bank冲突次数 | 0 |
| Sector Size | 每次内存事务的大小 | 32字节或128字节 |
举个例子，你在Nsight Compute里看到Global Load Efficiency只有30%，那就说明你的内存访问模式有问题——大概率是跨步访问或者不对齐。点开Source页，看看是哪一行代码导致的，然后针对性优化。
实操建议：每次写完kernel，先跑一遍profile。重点关注「Memory Workload Analysis」部分的「Global Load/Store Efficiency」。如果低于80%，就得回头检查代码了。
我记得有一次，帮一个团队优化推荐系统的embedding lookup。他们用nvprof一看，Global Load Efficiency只有12%。我让他们把embedding表按访问频率重新排列，让同一个warp内的线程访问连续地址。改完之后，效率升到92%，整体推理速度提升了3倍。嗯，这就是工具的价值。
知识体系总览
下面这张图，是我自己总结的内存访问模式优化框架。你可以把它当作一个检查清单：
这张图把今天讲的内容串起来了。从内存访问模式出发，分四个方向分析——合并访问、对齐、跨步访问、Bank Conflict。然后通过优化手段和Profiling工具，最终盯住那几个关键指标。你每次优化kernel，都可以拿这张图当checklist。
好了，今天的内容就到这里。记住：内存访问模式是GPU优化的第一道坎，跨过去，你的性能就能上一个台阶。