内存访问优化：全局内存合并访问、共享内存bank conflict、寄存器溢出与本地内存

各位同学，今天我们来聊聊内存访问优化。说实话，在GPU编程里，内存访问往往是性能瓶颈的“头号嫌疑人”。我见过太多项目，计算逻辑写得漂亮，结果一跑起来，带宽利用率低得可怜。说白了，GPU的算力再强，喂不饱数据也是白搭。

这一章，我会把内存访问的几个核心问题拆开揉碎了讲。全局内存怎么合并访问？共享内存的bank conflict怎么避免？寄存器溢出了怎么办？本地内存又是什么鬼？嗯，咱们一个一个来。

核心观点：GPU性能优化的本质，就是让数据流动起来。内存访问模式决定了带宽利用率，而带宽利用率决定了你的程序能跑多快。

全局内存合并访问

先说说全局内存。GPU的全局内存带宽很大，但延迟也很高。怎么高效利用？关键就是“合并访问”。

什么叫合并访问？说白了，就是让同一个warp里的32个线程，访问连续的内存地址。这样GPU可以把这些访问合并成一次或几次大的内存事务，效率极高。

反过来，如果每个线程访问的地址是分散的，GPU就得为每个线程单独发起内存事务，带宽利用率直线下降。我在项目中遇到过这种情况：一个简单的矩阵转置，因为没注意合并访问，性能差了10倍不止。

来看个例子：

// 合并访问版本（高效）
__global__ void kernel_coalesced(float *A, float *B, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        B[idx] = A[idx] * 2.0f;  // 连续访问
    }
}

// 非合并访问版本（低效）
__global__ void kernel_non_coalesced(float *A, float *B, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        B[idx] = A[idx * 2] * 2.0f;  // 跨步访问，无法合并
    }
}

你想想看，第一个版本里，线程0访问地址0，线程1访问地址1，线程2访问地址2……32个线程的访问正好是连续的，完美合并。第二个版本里，每个线程跳过一个元素，访问模式变成了0, 2, 4, 6……这就没法合并了。

我的经验：写kernel时，先想想warp里的线程访问模式。如果发现访问地址是跨步的，试着调整数据布局或者线程索引映射。我曾经把一个跨步访问的kernel改成连续访问，性能从20GB/s飙到了120GB/s。

共享内存bank conflict

共享内存是GPU上的“高速缓存”，延迟低、带宽高。但它有个坑——bank conflict。

共享内存被分成32个bank（以计算能力3.5+为例），每个bank的宽度是4字节。如果同一个warp里的多个线程同时访问同一个bank的不同地址，就会发生bank conflict，导致访问串行化。

举个例子：

// 有bank conflict的版本
__shared__ float s_data[32][32];
int tid = threadIdx.x;
int val = s_data[tid][tid];  // 所有线程访问同一列，bank冲突严重

// 无bank conflict的版本
__shared__ float s_data[32][32];
int tid = threadIdx.x;
int val = s_data[tid][0];  // 所有线程访问同一行，广播模式，无冲突

为什么会这样？因为共享内存的地址到bank的映射是：bank_index = (byte_address / 4) % 32。在第一个版本里，s_data[tid][tid]的地址是tid * 32 * 4 + tid * 4，算下来每个线程的bank索引都不一样，但如果有两个线程的bank索引相同，就冲突了。

实际上，s_data[tid][tid]这个模式，每个线程访问的是不同的行和不同的列，但地址计算后，你会发现线程0和线程16访问的是同一个bank（因为0 * 32 + 0 = 0，16 * 32 + 16 = 528，528 / 4 % 32 = 0）。所以，这是个典型的bank conflict场景。

避坑指南：我曾经在写卷积kernel时，没注意共享内存的bank映射，结果性能比预期差了一半。后来用padding技巧——在共享内存数组的每一行末尾加一个空元素——才解决了问题。比如__shared__ float s_data[32][33]，这样同一列的元素就分散到不同bank了。

寄存器溢出与本地内存

寄存器是GPU上最快的存储，但数量有限。每个线程能用的寄存器数量是固定的（比如256个），如果你在kernel里声明了太多局部变量，编译器就会把多余的变量“溢出”到本地内存。

本地内存听起来像“本地”，但实际上它用的是全局内存的同一套硬件，延迟高、带宽低。一旦发生寄存器溢出，性能会断崖式下跌。

来看个例子：

// 寄存器使用过多，可能溢出
__global__ void kernel_spill() {
    float a1, a2, a3, a4, a5, a6, a7, a8;
    float b1, b2, b3, b4, b5, b6, b7, b8;
    // ... 更多变量
    // 计算逻辑
}

// 优化后：减少寄存器使用
__global__ void kernel_no_spill() {
    float a[8];  // 使用数组，编译器可能优化
    // 或者分块计算，复用寄存器
}

怎么判断有没有寄存器溢出？用--ptxas-options=-v编译选项，看输出里的lmem字段。如果lmem大于0，说明有溢出。

我的习惯：写kernel时，我会刻意控制每个线程的寄存器使用量。如果发现lmem不为0，我会尝试：1) 减少局部变量数量；2) 把一些变量改成数组让编译器优化；3) 调整__launch_bounds__限制最大线程数，给每个线程更多寄存器。嗯，这些方法我都试过，效果不错。

知识体系总览

下面这张图，我把本章的核心知识点串起来了。你可以看到，内存访问优化的三个维度——全局内存、共享内存、寄存器/本地内存——是相互关联的。优化一个地方，可能影响另一个地方。比如减少寄存器使用，可能增加全局内存访问；增加共享内存使用，可能减少全局内存带宽压力。

实战建议

好了，理论讲完了，咱们来点实际的。我总结了几条优化建议，你在写kernel时可以对照着检查：

先看全局内存访问模式——确保warp内线程访问连续地址。如果不行，考虑转置数据或者调整线程索引。
再看共享内存bank冲突——用padding或者改变数据布局来避免。记住，s_data[32][33]比s_data[32][32]更安全。
最后检查寄存器使用——用-v编译选项看lmem。如果溢出，减少变量或调整__launch_bounds__。

一句话总结：内存访问优化的本质，就是让GPU的存储层次各司其职——全局内存要合并，共享内存要无冲突，寄存器要够用。这三件事做好了，你的kernel性能不会差。

嗯，这一章的内容就到这里。记住，优化不是一蹴而就的，需要反复调试和测量。我每次写新kernel，都会先跑一遍profiler，看看内存访问模式，再针对性优化。你也试试看？

📚 课程目录

内存访问优化：全局内存合并访问、共享内存bank conflict、寄存器溢出与本地内存

全局内存合并访问

共享内存bank conflict

寄存器溢出与本地内存

知识体系总览

实战建议