CUDA共享内存填充后仍存Bank冲突：为何填充31可完全消除？

按照CUDA矩阵转置的优化技巧，我测试了以下CUDA代码并得到对应的NCU性能分析结果。Bank冲突虽显著减少，但仍有残留。

// 存在Bank冲突
__global__ void setRowReadCol(int *out){
    __shared__ int tile[BDIMY][BDIMX];
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = idx;
    __syncthreads();
    out[idx] = tile[threadIdx.x][threadIdx.y];
}
// 预期无Bank冲突（填充1个元素）
__global__ void setRowReadColPad(int *out){
    __shared__ int tile[BDIMY][BDIMX + 1];  // BDIMX=BDIMY=32
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = idx;
    __syncthreads();
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

NCU性能分析结果 #

setRowReadCol(int*),  Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                                                          994
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                                                            0
    ---------------------------------------------------------------------- --------------- ------------------------------

  setRowReadColPad(int*),  Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                                                            2
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                                                            0
    ---------------------------------------------------------------------- --------------- ------------------------------

可见仍存在2次冲突事务。有趣的是，当填充大小调整为31时，Bank冲突被完全消除：

setRowReadColPad31(int*), Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                                                            0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                                                            0
    ---------------------------------------------------------------------- --------------- ------------------------------

请问有人能解释这一现象吗？

完整测试代码 #

#include <cuda_runtime.h>
#include <stdio.h>
#include <cstdlib>
#include <chrono>
#include <ctime>
#include <iostream>
#include <iomanip>
#include <stdio.h>
#include <stdarg.h>

#define CUDA_CHECK(call)                                      \
    do {                                                      \
        cudaError_t err = call;                               \
        if (err != cudaSuccess) {                             \
            std::cerr << "CUDA error: " << cudaGetErrorString(err) \
                      << " at " << __FILE__ << ":" << __LINE__ \
                      << std::endl;                           \
            exit(EXIT_FAILURE);                               \
        }                                                     \
    } while (0)


#define BDIMX 32
#define BDIMY 32

// 无冲突
__global__ void setRowReadRow (int *out){
    __shared__ int tile[BDIMY][BDIMX];
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = idx;
    __syncthreads();
    out[idx] = tile[threadIdx.y][threadIdx.x] ;
}

// 存在冲突
__global__ void setRowReadCol(int *out){
    __shared__ int tile[BDIMY][BDIMX];
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = idx;
    __syncthreads();
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

// 预期无冲突？
__global__ void setRowReadColPad(int *out){
    __shared__ int tile[BDIMY][BDIMX + 1];
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = idx;
    __syncthreads();
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

__global__ void setRowReadColPad31(int *out) {
    __shared__ int tile[BDIMY][BDIMX + 31];
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = idx; __syncthreads();
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

int main(int argc, char **argv)
{
    // 初始化设备
    int dev = 0;
    cudaDeviceProp deviceProp;
    CUDA_CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("%s at ", argv[0]);
    printf("device %d: %s ", dev, deviceProp.name);
    CUDA_CHECK(cudaSetDevice(dev));

    cudaSharedMemConfig pConfig;
    CUDA_CHECK(cudaDeviceGetSharedMemConfig ( &pConfig ));
    printf("with Bank Mode:%s ", pConfig == 1 ? "4-Byte" : "8-Byte");

    // 设置数组大小
    int nx = BDIMX;
    int ny = BDIMY;

    bool iprintf = 0;

    if (argc > 1) iprintf = atoi(argv[1]);

    size_t nBytes = nx * ny * sizeof(int);

    // 执行配置
    dim3 block (BDIMX, BDIMY);
    dim3 grid  (1, 1);
    printf("<<< grid (%d,%d) block (%d,%d)>>>\n", grid.x, grid.y, block.x,
           block.y);

    // 分配设备内存
    int *d_C;
    CUDA_CHECK(cudaMalloc((int**)&d_C, nBytes));
    int *gpuRef  = (int *)malloc(nBytes);

    CUDA_CHECK(cudaMemset(d_C, 0, nBytes));
    setRowReadRow<<<grid, block>>>(d_C);
    CUDA_CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    CUDA_CHECK(cudaMemset(d_C, 0, nBytes));
    setRowReadCol<<<grid, block>>>(d_C);
    CUDA_CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    CUDA_CHECK(cudaMemset(d_C, 0, nBytes));
    setRowReadColPad<<<grid, block>>>(d_C);
    CUDA_CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    CUDA_CHECK(cudaMemset(d_C, 0, nBytes));
    setRowReadColPad31<<<grid, block>>>(d_C);
    CUDA_CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    CUDA_CHECK(cudaFree(d_C));
    free(gpuRef);
    return EXIT_SUCCESS;
}

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