CUDA共享内存二维数组的Bank Conflict问题排查与优化咨询

问题原因分析 #

你的共享内存二维数组采用行优先存储，ms[N][32]这类数组中，同一列的所有元素会被映射到同一个共享内存bank：

• 每个行有32个float（刚好对应CUDA默认的32个共享内存bank），ms[i][threadIdx.x]的内存word索引为i*32 + threadIdx.x，对32取模后结果恒等于threadIdx.x，也就是说ms[0][t]、ms[1][t]……ms[9][t]全落在bankt上。
• 当32线程的warp执行循环时，每个线程t会连续10次向同一个bankt发起写请求（对应循环的10次迭代）。共享内存每个bank每个周期只能处理一个请求，这10次请求需要串行化执行，因此被Nsight Compute的l1_conflicts_shared_nway指标统计为10路冲突。
• 另外你声明的ms、is、ds三个同维度数组在共享内存中连续分配，导致is[i][t]和ds[i][t]也会落在bankt上，进一步增加了同一bank的访问压力，但核心冲突来源是单线程对同一bank的连续10次访问。

优化建议 #

• 数组填充（Padding）：修改数组列数，让每行元素个数不是32的整数倍，打破同一列元素的bank映射。例如将数组定义为：

  #define N 10
  __shared__ float ms[N][33];  // 增加1个填充列
  __shared__ float is[N][33];
  __shared__ float ds[N][33];
  

  此时ms[i][t]的word索引为i*33 + t，bank索引变为(i + t) %32，不同行的同一列元素会分散到不同bank，彻底避免单线程连续访问同一bank。后续访问时只需使用0-31列即可。

• 调整访问模式：如果初始化阶段允许，改为按行访问。例如调整数组维度为ms[32][N]，让线程t负责初始化ms[t][0..9]：

  for (int i = 0; i < N; ++i) {
      ms[threadIdx.x][i] = 0.0f;
      is[threadIdx.x][i] = 0.0f;
      ds[threadIdx.x][i] = 0.0f;
  }
  

  这种模式下每个线程访问的元素会分布在不同bank，无冲突。

• 修改共享内存bank大小：如果你的CUDA设备支持，可通过运行时API将bank大小设置为8字节：

  cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
  

  此时每个bank可存储两个float，同一列的元素会分散到不同bank，但这种方法可能影响其他代码的内存访问模式，需谨慎测试。

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