CUDA性能调优实战：用Padding技巧解决共享内存Bank Conflict，让矩阵转置快3倍

CUDA性能调优实战：用Padding技巧解决共享内存Bank Conflict，让矩阵转置快3倍

在GPU并行计算的世界里，共享内存就像是一把双刃剑——用得好能带来数量级的性能提升，用得不好反而会成为程序瓶颈。最近在处理一个大型矩阵转置任务时，我发现了一个有趣的现象：明明已经使用了共享内存，但性能却比预期低了近40%。通过Nsight Compute工具分析，罪魁祸首竟然是Bank Conflict这个隐藏的性能杀手。

1. 揭开Bank Conflict的神秘面纱

现代GPU的共享内存被划分为32个独立的存储体（bank），每个bank可以同时响应一个内存请求。当多个线程试图访问同一个bank的不同地址时，这些访问会从并行变为串行——这就是Bank Conflict的本质。

以一个简单的矩阵转置为例：

__global__ void transposeNaive(float *out, float *in, int nx, int ny) {
    int ix = blockIdx.x * blockDim.x + threadIdx.x;
    int iy = blockIdx.y * blockDim.y + threadIdx.y;
    if (ix < nx && iy < ny) {
        out[ix * ny + iy] = in[iy * nx + ix]; // 列主序写入
    }
}

这个朴素实现存在明显的合并访问问题。改用共享内存后：

__global__ void transposeSmem(float *out, float *in, int nx, int ny) {
    __shared__ float tile[BDIMY][BDIMX];
    // ... 数据加载到共享内存 ...
    tile[threadIdx.y][threadIdx.x] = in[iy * nx