内存访问模式：合并访问与银行冲突
各位同学，今天我们来聊聊显存优化里最核心的两个概念——合并访问和银行冲突。这两个东西，说白了就是决定你的GPU到底能跑多快的关键。我刚开始学CUDA的时候，总觉得显存带宽那么大，随便读写不就行了？结果一跑性能测试，直接被打脸。嗯，今天我们就来彻底搞明白它们。
一、合并访问（Coalesced Access）原理
先问大家一个问题：GPU的显存带宽很高，但为什么你的核函数经常跑不满带宽？
答案就在合并访问上。GPU在执行内存访问时，是以warp（32个线程）为单位的。当一个warp里的所有线程访问连续的、对齐的内存地址时，硬件会把多次小请求合并成一次大请求。这就是合并访问。
核心要点：合并访问能让一次内存事务传输更多数据，从而充分利用显存带宽。不合并的话，每次只传一点点，带宽利用率极低。
举个例子，假设每个线程要读一个float（4字节），32个线程就是128字节。如果地址连续且对齐，一次事务就能搞定。如果地址是乱序的，那就得多次事务，性能直接打折扣。
我个人习惯在写核函数时，先看一眼全局内存的访问模式。如果发现是跨步访问，我会优先考虑能不能转置数据或者调整线程布局。
二、对齐与步长
合并访问有两个关键条件：对齐和连续。
- 对齐：起始地址必须是32字节、64字节或128字节的倍数。这取决于你的数据类型和硬件代际。
- 连续：线程0访问地址A，线程1访问A+1，线程2访问A+2……以此类推。
但实际项目中，我们经常遇到步长访问。比如每个线程跳过N个元素再读下一个。这时候合并访问就失效了。
避坑指南：我曾经在一个图像处理项目里，为了图方便，直接用线程ID乘以一个步长去读像素数据。结果性能惨不忍睹，只有理论带宽的10%。后来改成连续读取，再在共享内存里做重排，性能直接翻了6倍。
来看一个对比代码：
// ❌ 非合并访问：步长为2
__global__ void stride_access(float *in, float *out, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
out[idx] = in[idx * 2]; // 每个线程跳过一个元素
}
}
// ✅ 合并访问：连续读取
__global__ void coalesced_access(float *in, float *out, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
out[idx] = in[idx]; // 连续地址
}
}你想想看，第一个核函数里，线程0读地址0，线程1读地址2，线程2读地址4……这中间的空隙全浪费了。第二个核函数才是正确的打开方式。
三、银行冲突（Bank Conflict）详解
说完全局内存，咱们来看看共享内存。共享内存虽然快，但它有个坑——银行冲突。
共享内存被划分成32个银行（bank），每个bank宽度为4字节。每个bank在每个时钟周期只能响应一次访问。如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致访问串行化。
| 访问模式 | 是否冲突 | 性能影响 |
|---|---|---|
| 所有线程访问同一地址（广播） | 无冲突 | 一次事务 |
| 线程访问不同bank | 无冲突 | 一次事务 |
| 多个线程访问同一bank的不同地址 | 有冲突 | 多次事务 |
举个例子，共享内存地址0、4、8、12……属于bank0。如果线程0读地址0，线程1读地址4，线程2读地址8……它们全挤在bank0上，这就炸了。
我的经验：在写矩阵转置或者卷积核函数时，我经常在共享内存里加一个padding列。比如原本声明__shared__ float s[32][32]，改成__shared__ float s[32][33]。这样原本在同一bank的地址就被错开了，冲突就没了。这招我用了好多年，屡试不爽。
为什么会这样？因为加了一列padding后，原本在同一行的相邻元素，它们的地址偏移变成了33个float，也就是132字节。132除以4等于33，33对32取模是1，所以它们落到了不同的bank上。
四、知识体系结构图
下面我用一张SVG图来总结本章的核心逻辑，方便你理解整个脉络：
五、实战建议
最后，我给大家几个实战中可以直接用的建议：
- 先分析再动手：用
nvprof或Nsight Compute看一下你的内存访问模式。如果发现Global Load Efficiency低于80%，那基本就是合并访问没做好。 - 共享内存padding：遇到银行冲突时，先试试加一列padding。这招简单粗暴，但效果立竿见影。
- 数据布局先行：很多时候，合并访问的问题出在数据布局上。比如结构体数组（AoS）改成数组结构体（SoA），访问模式就自动变连续了。
记住：显存优化的本质，就是让硬件用最少的次数搬最多的数据。合并访问和银行冲突，一个管全局内存，一个管共享内存，两者都优化好了，你的核函数才能跑出接近理论峰值的性能。
好了，这一章的内容就到这里。希望大家在实际项目中多动手、多分析，慢慢就会形成直觉——看到一段代码，就能大概猜到它的内存访问效率如何。