内存访问优化（上）：合并访问原理与实现

各位同学，今天我们来聊聊CUDA优化里最核心的一个话题——内存访问。说实话，我见过太多新手写的kernel，计算逻辑没问题，但跑起来就是慢。一分析，问题全出在内存访问上。GPU的算力其实很猛，但如果你喂数据的方式不对，它就只能干等着。

为什么内存访问这么重要？

先看一组数据。我随便拿一块A100来说，它的计算吞吐量能达到312 TFLOPS（FP16），但全局内存带宽只有2 TB/s左右。你算算，如果每个浮点运算都需要从显存读数据，那计算单元大部分时间都在空转。说白了，内存访问就是GPU性能的咽喉。

我在优化一个3D点云处理项目时遇到过这种情况。同样的算法，我改了一下数据排布方式，性能直接提升了8倍。当时团队里的人都惊呆了。其实没什么神奇的，就是让内存访问变得"合并"了而已。

合并访问（Coalesced Access）原理

什么是合并访问？简单说，就是一个warp里的32个线程，访问连续的内存地址。GPU的内存控制器会把这32次访问合并成一次或几次大的内存事务。你想想看，一次读128字节和32次读4字节，效率能一样吗？

为什么会这样？因为GPU的内存系统设计就是按"合并访问"来优化的。每个内存事务可以传输32字节、64字节或128字节。如果warp里的线程访问地址是连续的，一个128字节的事务就能满足32个线程的需求（每个线程读4字节）。

我画了一张图，帮你理解这个逻辑：

Global Memory带宽利用率计算

带宽利用率，说白了就是实际用到的带宽除以理论峰值带宽。我习惯用这个公式：

利用率 = (实际读取字节数) / (理论峰值带宽 × 执行时间)

但这里有个坑。你实际从显存读了多少字节，和你的代码里写了多少读操作，不是一回事。因为GPU会做内存事务合并。

举个例子：

访问模式	线程请求字节	实际传输字节	利用率
完全合并（连续）	32 × 4 = 128B	128B（1次事务）	100%
跨步访问（步长2）	32 × 4 = 128B	256B（2次事务）	50%
随机访问	32 × 4 = 128B	可能1024B+	< 12.5%

看到没？同样的128字节请求，实际传输量可以差8倍以上。我在优化一个矩阵转置的kernel时，就踩过这个坑。一开始写的版本带宽利用率只有15%，改完合并访问后直接飙到85%。

对齐访问（Aligned Access）

对齐访问和合并访问是两兄弟。合并是说线程间地址连续，对齐是说起始地址是32/64/128的倍数。

为什么对齐重要？因为GPU的内存控制器是按"段"来工作的。每个段128字节，起始地址必须是128的倍数。如果你的数据起始地址是100，那第一个段从0开始，第二个段从128开始。你的数据横跨两个段，就得发两次事务。

跨步访问（Strided Access）

跨步访问是合并访问的"天敌"。什么叫跨步？就是线程i访问地址A + i × stride，其中stride > 1。

最常见的场景是二维数组按列访问。比如一个N×N的矩阵，按行存储。线程i要访问第i行第j列，那地址就是base + i × N + j。如果j固定，i变化，那步长就是N。N越大，跨步越大，合并效果越差。

// ❌ 跨步访问（按列访问）
__global__ void col_access(float *mat, float *out, int N) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    // 每个线程访问不同行、同一列
    out[row * N + col] = mat[row * N + col];  // 如果col固定，row变化，步长为N
}

// ✅ 合并访问（按行访问）
__global__ void row_access(float *mat, float *out, int N) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    // 每个线程访问同一行、不同列
    out[row * N + col] = mat[row * N + col];  // 如果row固定，col变化，步长为1
}

你看，同样的数据，只是交换了row和col的角色，访问模式就完全不同了。我在优化一个图像处理算法时，就是靠这个技巧，把处理时间从12ms降到了2ms。

实战：如何写出合并访问的代码

我总结了几个实用原则：

1. 让threadIdx.x对应最内层维度。因为threadIdx.x是变化最快的线程索引，让它对应连续的内存地址。
2. 使用一维布局访问二维数据。比如用blockIdx.x * blockDim.x + threadIdx.x作为全局索引，然后手动计算行和列。
3. 考虑数据转置。如果算法必须按列访问，可以先把矩阵转置，再按行访问。
4. 使用共享内存做中转。先按合并方式读到共享内存，再在共享内存里做非合并访问。共享内存没有合并访问的限制。

如何检测合并访问问题

我常用的工具是NVIDIA Nsight Compute。它有个指标叫Global Load Efficiency和Global Store Efficiency。如果这个值低于80%，你的内存访问就有问题。

另外，你也可以用CUDA自带的cudaMemGetInfo来监控显存带宽使用情况。但说实话，最直接的方法还是看kernel的执行时间。改一下访问模式，时间降了，那就是合并访问的功劳。

嗯，今天的内容就到这里。合并访问是CUDA优化的基本功，但很多人容易忽略。你想想看，一个简单的数据排布调整，就能带来几倍的性能提升，这性价比多高啊。

记住：让相邻线程访问相邻地址。这句话值千金。

想要更多资料，联系可可