如何消除CUDA核函数内存Bank冲突？改写后的核函数是否正确？

我需要解决CUDA核函数中的内存Bank冲突问题，先给出一段存在该冲突的原核函数代码，之后自行编写了改进版（BLOCK_SIZE设为32），想确认这个改进方式是否符合规范。

原核函数代码（存在Bank冲突） #

__global__ void matrixOuterProductKernel(float* A, float* B, float* C, int K) {
    // 计算当前线程的行、列索引
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 分配共享内存存储相关向量和结果片段
    __shared__ float sA[K][K];
    __shared__ float sB[K][K];
    __shared__ float sC[K][K];
    
    // 初始化结果片段为0
    sC[threadIdx.y][threadIdx.x] = 0.0;
    
    // 循环遍历向量并累加结果片段
    for (int i = 0; i < K; i++) {
        // 将向量的相关元素加载到共享内存
        sA[threadIdx.y][i] = A[row * K + i];
        sB[i][threadIdx.x] = B[i * K + col];
        
        // 等待所有线程完成向量加载
        __syncthreads();
        
        // 计算外积并累加到结果片段
        for (int j = 0; j < K; j++) {
            sC[threadIdx.y][threadIdx.x] += sA[threadIdx.y][j] * sB[j][threadIdx.x];
        }
        
        // 等待所有线程完成当前迭代，再加载下一组向量
        __syncthreads();
    }
    
    // 将结果片段写回全局内存
    C[row * K + col] = sC[threadIdx.y][threadIdx.x];
}

改进后的核函数代码（BLOCK_SIZE=32） #

__global__ void matrixOuterProductKernel(const double* a, const double* b, double* result, const int n) {
    __shared__ double sA[BLOCK_SIZE * BLOCK_SIZE];
    __shared__ double sB[BLOCK_SIZE * BLOCK_SIZE];

    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;

    int row = threadIdx.y;
    int col = threadIdx.x;

    int globalRow = blockRow * blockDim.y + row;
    int globalCol = blockCol * blockDim.x + col;

    double sum = 0.0;

    for (int i = 0; i < (n + BLOCK_SIZE - 1) / BLOCK_SIZE; i++) {
        int aRow = globalRow * n + i * BLOCK_SIZE + col;
        int bCol = (i * BLOCK_SIZE + row) * n + globalCol;

        if (globalRow < n && (i * BLOCK_SIZE + col) < n) {
            sA[row * BLOCK_SIZE + col] = a[aRow];
        } else {
            sA[row * BLOCK_SIZE + col] = 0.0;
        }

        if ((i * BLOCK_SIZE + row) < n && globalCol < n) {
            sB[row * BLOCK_SIZE + col] = b[bCol];
        } else {
            sB[row * BLOCK_SIZE + col] = 0.0;
        }

        __syncthreads();

        for (int j = 0; j < BLOCK_SIZE; j++) {
            sum += sA[row * BLOCK_SIZE + j] * sB[j * BLOCK_SIZE + col];
        }

        __syncthreads();
    }

    if (globalRow < n && globalCol < n) {
        result[globalRow * n + globalCol] = sum;
    }
}

改进方案的规范性验证 #

原代码的Bank冲突根源 #

原代码使用二维共享内存数组sA[K][K]、sB[K][K]，当K=32（与CUDA默认的32个内存Bank数量一致）时，同一Warp的线程访问sA[threadIdx.y][i]时，共享内存地址偏移为threadIdx.y*K + i，所有线程的地址低5位（Bank索引位）均为i，导致所有线程同时访问同一个Bank，触发严重的Bank冲突，大幅降低内存访问效率。

改进方案的合规性分析 #

你的改进方案完全符合CUDA优化的最佳实践，有效解决了Bank冲突问题，具体体现在：

• 分块（Tiling）策略：将大矩阵拆分为BLOCK_SIZE*BLOCK_SIZE的块处理，减少全局内存访问次数，这是CUDA矩阵运算优化的标准做法。
• 共享内存访问模式优化：
  • 加载sA时，同一Warp内的线程按row*BLOCK_SIZE + col的索引存储，地址低5位对应col（0~31），恰好映射到不同的内存Bank，无冲突。
  • 计算阶段访问sA[row*BLOCK_SIZE + j]和sB[j*BLOCK_SIZE + col]时，同一Warp内线程的地址低5位分别为j和col，均覆盖0~31的Bank索引，避免冲突。
• 边界合法性检查：针对矩阵大小非BLOCK_SIZE整数倍的情况，添加越界判断并赋值0，避免非法内存访问，保障核函数的鲁棒性。
• 寄存器替代共享内存存储中间结果：用寄存器变量sum代替原代码的共享内存sC存储累加结果，减少共享内存占用和访问开销，进一步提升性能。

综上，这个改进方式是规范且有效的，完全解决了原核函数中的内存Bank冲突问题。

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