访存优化：合并访问、共享内存与数据预取

大家好，欢迎来到《算子优化与显存管理实战》的第一章。

咱们今天聊的，是GPU优化的基本功——访存优化。说白了，就是怎么让数据在GPU里跑得更快。我刚开始接触CUDA时，觉得只要把计算扔到GPU上就完事了，结果一跑，性能还不如CPU。后来才发现，瓶颈根本不在计算，而在访存。

你想想看，GPU的计算单元多到爆炸，但显存带宽就那么点。如果数据喂不饱计算单元，再强的算力也是白搭。所以，访存优化是绕不开的第一关。

1. 合并访问（Coalesced Access）

先讲合并访问。这个概念，我当年理解了好久。

GPU的显存是分片的，每个内存控制器管一片。当线程束（warp）里的32个线程同时访问显存时，硬件会尽量把这些访问请求合并成一次或少数几次大的事务。如果访问地址是连续的，那就完美合并，一次搞定。如果地址是乱序的，那就得拆成很多次小事务，带宽利用率直线下降。

我举个例子。假设你有一个float数组，每个线程读一个元素：

// 合并访问：连续地址
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 线程0读data[0]，线程1读data[1]... 完美合并

// 非合并访问：跨步访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx * 2];  // 线程0读data[0]，线程1读data[2]... 地址不连续

第一种写法，32个线程读的是连续地址，硬件一次搞定。第二种写法，地址间隔了2个float，硬件得拆成两次事务，带宽利用率直接腰斩。

我在项目中遇到过一个问题：一个矩阵转置的算子，性能死活上不去。后来一查，发现是列方向读取时，线程访问的地址跨步太大，完全没有合并。改成按行读取、按列写入后，性能提升了3倍。嗯，就是这么夸张。

2. 共享内存的使用与bank conflict

共享内存是GPU里的一块特殊存储，速度比全局显存快得多，但容量小（通常几十KB）。它的作用，说白了就是给线程块内的线程提供一个快速交换数据的空间。

但共享内存有个坑——bank conflict。

共享内存被分成32个bank，每个bank每周期只能响应一个请求。如果多个线程同时访问同一个bank的不同地址，就会发生bank conflict，导致访问串行化。如果所有线程访问同一个地址，反而没事（广播）。

我画个图帮你理解：

为什么会这样？因为地址到bank的映射是固定的：地址 % 32 决定它落在哪个bank。如果你访问的地址间隔正好是32的倍数，那所有线程都会打到同一个bank上，性能直接崩掉。

我个人的习惯是，在写共享内存代码时，先算一下访问模式。比如做矩阵分块时，给共享内存数组多分配一列（padding），就能避免bank conflict：

// 有bank conflict的写法
__shared__ float s_data[32][32];
float val = s_data[threadIdx.y][threadIdx.x];  // 同一行，不同列，无冲突
float val2 = s_data[threadIdx.x][threadIdx.y]; // 同一列，不同行，有冲突！

// 加padding避免冲突
__shared__ float s_data[32][33];  // 多一列
float val2 = s_data[threadIdx.x][threadIdx.y]; // 现在不冲突了

我曾经在一个卷积算子的优化中，因为没注意bank conflict，性能一直上不去。后来加了padding，吞吐量直接提升了40%。嗯，细节决定成败。

3. 数据预取（Prefetching）

数据预取，说白了就是提前把数据搬到离计算单元更近的地方。GPU的显存层次是：全局显存 → L2缓存 → L1缓存/共享内存 → 寄存器。每一层速度越来越快，但容量越来越小。

预取的思路很简单：在当前数据计算的同时，把下一批数据从慢速存储搬到快速存储。这样计算单元就不用等数据了。

我举个例子，在矩阵乘法中，我们经常用分块+预取：

// 伪代码：矩阵乘法中的双缓冲预取
__shared__ float s_A[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float s_B[BLOCK_SIZE][BLOCK_SIZE];

// 预取第一块
s_A[threadIdx.y][threadIdx.x] = A[blockIdx.y * BLOCK_SIZE + threadIdx.y][0 + threadIdx.x];
s_B[threadIdx.y][threadIdx.x] = B[0 + threadIdx.y][blockIdx.x * BLOCK_SIZE + threadIdx.x];
__syncthreads();

for (int k = 0; k < K; k += BLOCK_SIZE) {
    // 预取下一块（异步）
    if (k + BLOCK_SIZE < K) {
        s_A[threadIdx.y][threadIdx.x] = A[blockIdx.y * BLOCK_SIZE + threadIdx.y][k + BLOCK_SIZE + threadIdx.x];
        s_B[threadIdx.y][threadIdx.x] = B[k + BLOCK_SIZE + threadIdx.y][blockIdx.x * BLOCK_SIZE + threadIdx.x];
    }
    
    // 计算当前块
    for (int i = 0; i < BLOCK_SIZE; i++) {
        sum += s_A[threadIdx.y][i] * s_B[i][threadIdx.x];
    }
    
    __syncthreads();  // 确保预取完成
}

你看，这里用共享内存做缓冲，计算和预取重叠，带宽利用率就上去了。

我记得有一次优化一个稀疏矩阵算子，想用预取，结果发现数据访问完全随机，预取命中率不到10%，性能反而下降了。后来改成按需加载，配合合并访问，效果才好。

总结一下

访存优化，核心就三件事：

• 合并访问：让warp内的线程访问连续地址，这是基础。
• 共享内存与bank conflict：用好共享内存，避开bank冲突，这是进阶。
• 数据预取：计算和访存重叠，这是高手玩法。

这三层做好了，你的算子性能基本不会差。我见过太多人一上来就调计算指令，结果发现瓶颈在访存，白费功夫。所以，先搞定访存，再谈计算优化。

好了，第一章就到这里。记住，GPU优化的本质，就是让数据流动起来，别让计算单元闲着。