为何仅对transposedTile共享内存填充而不对aTile进行填充？

• 使用全局内存合并读取优化处理跨步访问

上述链接指出：

__global__ void coalescedMultiply(float *a, float *c, int M)
{
  __shared__ float aTile[TILE_DIM][TILE_DIM],
                   transposedTile[TILE_DIM][TILE_DIM];
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
  transposedTile[threadIdx.x][threadIdx.y] =
      a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
      threadIdx.x];  
  __syncthreads();
  for (int i = 0; i < TILE_DIM; i++) {
      sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
  }
  c[row*M+col] = sum;
}

(...)
此类多路Bank冲突开销极大。简单的解决方法是对共享内存数组进行填充，增加一列，如下代码所示：

__shared__ float transposedTile[TILE_DIM][TILE_DIM+1];

问题 #

为何仅对右侧的transposedTile共享内存进行填充，而非同时对transposedTile和aTile两者进行填充？

解答 #

核心原因是两个共享内存数组的访问模式是否会触发严重的Bank冲突：

1. aTile的访问无严重Bank冲突 #

aTile采用行优先存储，其读写模式都不会触发高开销的多路Bank冲突：

• 写入阶段：同一warp内的线程按threadIdx.x递增访问aTile[threadIdx.y][threadIdx.x]，对应共享内存的连续地址，每个线程访问不同的Bank，完全无冲突。
• 读取阶段：循环中线程访问aTile[threadIdx.y][i]，即使同一warp内的线程同时访问不同行的同一列，这种访问的冲突程度极低，不会产生显著性能开销，无需额外填充。

2. transposedTile的写入会触发严重Bank冲突 #

transposedTile是转置存储，写入transposedTile[threadIdx.x][threadIdx.y]时，同一warp内的线程（threadIdx.y相同，threadIdx.x从0到31）会同时访问共享内存的同一个Bank（当TILE_DIM为32的倍数时），触发32路Bank冲突，每个内存访问需要多个周期才能完成，开销极大。

给transposedTile增加一列填充后，其内存布局变为TILE_DIM][TILE_DIM+1，此时transposedTile[x][y]的内存地址为x*(TILE_DIM+1) + y，同一warp线程的访问地址对32取模结果不再一致，彻底避免了Bank冲突。

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