CUDA中共享内存一维索引较二维索引读取速度慢的原因探究

核心问题：共享内存银行冲突 #

CUDA共享内存被划分为32个独立的内存银行（针对支持SM 2.0及以上的设备），每个银行可并行处理一次32位数据访问。当多个线程同时访问同一银行时，访问会被串行化，导致内存事务数激增、性能大幅下降。

「Good」代码的访问模式 #

代码中共享内存定义为：

__shared__ float tile[tile_dim][tile_dim + 1];

这里的tile_dim + 1是特意添加的填充位，目的是打破内存地址的对齐规律，避免银行冲突。

在「Good」行：

out[y * columns + x] = tile[threadIdx.x][threadIdx.y];

线程访问共享内存的地址计算为：
基地址 + threadIdx.x * (tile_dim + 1) * sizeof(float) + threadIdx.y * sizeof(float)
由于tile_dim + 1 = 33，相邻threadIdx.x对应的地址间隔为33*4=132字节，映射到银行编号时，不同线程的访问会均匀分布到32个银行中，无银行冲突，因此内存事务数最少，性能最优。

「Bad」代码的访问模式 #

「Bad」行强制将二维数组转为一维数组访问：

out[y * columns + x] = ((float*)tile)[threadIdx.x * tile_dim + threadIdx.y];

此时索引计算为threadIdx.x * 32 + threadIdx.y，对应的地址间隔为32*4=128字节，映射到银行编号时，同一threadIdx.y的所有线程（一个warp内的32个线程）都会访问同一个银行，引发严重的银行冲突。

这种情况下，每个warp的共享内存访问需要拆分为32次串行事务，直接导致共享内存加载事务数从20万暴增至6800万，运行时间几乎翻倍。

内容的提问来源于stack exchange，提问作者user1166