Code

// 有 Bank Conflict 的 Kernel
__global__ void kernelWithBankConflict(float *input, float *output) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    tile[threadIdx.x][threadIdx.y] = input[x * TILE_DIM + y];
    __syncthreads();
    output[x * TILE_DIM + y] = tile[threadIdx.x][threadIdx.y];
}

// 避免 Bank Conflict 的 Kernel
__global__ void kernelWithoutBankConflict(float *input, float *output) {
    __shared__ float tile[TILE_DIM][TILE_DIM + 1];
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    tile[threadIdx.x][threadIdx.y] = input[x * TILE_DIM + y];
    __syncthreads();
    output[x * TILE_DIM + y] = tile[threadIdx.x][threadIdx.y];
}

int main() {
    ...
    dim3 blockDim(TILE_DIM, TILE_DIM);
    dim3 gridDim(1, 1);
    // 执行 Kernel (有 Bank Conflict)
    kernelWithBankConflict<<<gridDim, blockDim>>>(d_input, d_output);

    // 执行 Kernel (避免 Bank Conflict)
    kernelWithoutBankConflict<<<gridDim, blockDim>>>(d_input, d_output);
    ...
}

input[x * TILE_DIM + y]; : x，y 表示每个线程自己的全局索引，x * TILE_DIM 表示目标位置所在的行，+ y 表示目标位置所在的列，所以行偏移后，列偏移，就得到了目标位置 index。

既然在同一个 GPU 上共享内存的 Bank 宽度是 4 字节，那么处理 float (4 字节) 和 double (8 字节) 数据类型时，需要考虑如何有效地访问共享内存，以减少 Bank conflict

通过 padding 构建新的数据类型：

struct DoubleData {
    double value;
    char padding[4]; // 添加 4 字节 padding，确保对齐
};

这里的 padding 并不是为了让 double 类型本身对齐（因为 double 本身通常会按照 8 字节对齐），而是为了避免不同线程访问共享内存时发生 Bank conflict。

上述，Bank 宽 4 bytes的组织方式，如果double数据不 padding，会导致 Bank conflict 如下：

Thread 0 访问 shared_double[0].value (Addr 0) -> Bank 0
Thread 1 访问 shared_double[1].value (Addr 8) -> Bank 2
Thread 2 访问 shared_double[2].value (Addr 16) -> Bank 4
...
Thread 7 访问 shared_double[7].value (Addr 56) -> Bank 14
Thread 8 访问 shared_double[8].value (Addr 64) -> Bank 16
...
Thread 15 访问 shared_double[15].value (Addr 120) -> Bank 30
Thread 16 访问 shared_double[16].value (Addr 128) -> Bank 0

padding 的主要作用是改变线程访问共享内存时，数据到 Bank 的映射关系。