CUDA矩阵乘法共享内存Bank冲突问题求助

问题分析 #

你的代码中共享内存采用连续的dim×dim布局，当矩阵尺寸超过128时触发Bank冲突，核心原因是共享内存的访问模式导致线程束内多个线程的访问请求映射到同一Bank。在Ampere架构（如A4000）的默认共享内存配置下，连续的行索引布局会让特定偏移的访问落在同一Bank，当矩阵尺寸增大后，这类冲突的频次被Nsight Compute清晰检测到。

解决方案：为共享内存添加Padding #

通过在共享内存数组的行末尾添加一个额外元素，打破连续布局的Bank映射规则，从根源避免冲突。具体修改如下：

修改核心Kernel代码 #

#define BLOCK_DIM 16
// 为每行添加1个padding元素，避免Bank冲突
#define SHMEM_SIZE ((BLOCK_DIM + 1) * BLOCK_DIM)

__global__ void matMulKernelSharedMemory(float *matrixA, float *matrixB, float *matrixC, int dimSize)
{
    __shared__ float s_matrixA[SHMEM_SIZE];
    __shared__ float s_matrixB[SHMEM_SIZE];

    int col = blockIdx.x * BLOCK_DIM + threadIdx.x;
    int row = blockIdx.y * BLOCK_DIM + threadIdx.y;

    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int dim = BLOCK_DIM;

    float value = 0.0f;
    // 标记当前线程是否在有效矩阵范围内
    bool valid = (row < dimSize) && (col < dimSize);

    for (int i = 0; i < ((dimSize + dim - 1) / dim); i++)
    {
        // 加载共享内存A，仅访问有效范围的全局内存
        if (row < dimSize && (i * dim + tx) < dimSize)
            s_matrixA[ty * (dim + 1) + tx] = matrixA[row * dimSize + i * dim + tx];
        else
            s_matrixA[ty * (dim + 1) + tx] = 0.0f;

        // 加载共享内存B，仅访问有效范围的全局内存
        if (col < dimSize && (i * dim + ty) < dimSize)
            s_matrixB[ty * (dim + 1) + tx] = matrixB[(i * dim + ty) * dimSize + col];
        else
            s_matrixB[ty * (dim + 1) + tx] = 0.0f;

        __syncthreads();

        // 计算矩阵乘积，使用带padding的共享内存索引
        for (int j = 0; j < dim; j++)
        {
            if (valid)
                value += s_matrixA[ty * (dim + 1) + j] * s_matrixB[j * (dim + 1) + tx];
        }
        __syncthreads();
    }

    if (valid)
        matrixC[row * dimSize + col] = value;
}

维度配置调整 #

保持block尺寸为16×16或32×32，对应修改BLOCK_DIM宏即可：

int BLOCK_DIM = 16;
int blocks = (dimSize + BLOCK_DIM - 1) / BLOCK_DIM;
dim3 dimGrid(blocks, blocks);
dim3 dimBlock(BLOCK_DIM, BLOCK_DIM);

关键改进说明 #

1. 共享内存Padding：通过ty*(dim+1)+tx的索引方式，让相邻行的元素落在不同Bank，彻底消除Bank冲突的可能。
2. 边界检查：添加矩阵范围判断，避免矩阵尺寸非block倍数时的越界访问，保证计算结果的正确性。
3. 灵活适配：通过宏定义BLOCK_DIM，可以轻松切换16×16或32×32的block配置，无需修改核心计算逻辑。

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