CUDA共享内存Bank冲突预期外时序问题：为何冲突与无冲突版耗时相近？

我尝试复现Bank冲突场景，针对一个warp（32线程）访问32个32位整数的两种场景开展基准测试：

• 无Bank冲突场景（offset=1）
• 存在Bank冲突场景（offset=32，所有线程访问bank 0）

核心kernel代码示例：

__global__ void kernel(int offset) {

    __shared__ uint32_t shared_memory[MEMORY_SIZE];

    // 初始化共享内存
    if (threadIdx.x == 0) {
        for (int i = 0; i < MEMORY_SIZE; i++) 
            shared_memory[i] = i;
    }

    __syncthreads();

    uint32_t index = threadIdx.x * offset;

    // 2048 / 32 = 64 
    for (int i = 0; i < 64; i++)
    {
        shared_memory[index] += index * 10;

        index += 32;
        index %= MEMORY_SIZE;

        __syncthreads();   
     }
}

我原本预期offset=32的冲突版本因访问序列化会运行更慢，但实际两者耗时相近，这是为何？

核心原因分析 #

你的测试代码存在几个关键问题，导致Bank冲突的性能损耗被完全掩盖：

1. 循环内同步操作的开销主导运行时间
   每次循环迭代都调用__syncthreads()，这个同步操作的延迟远高于共享内存访问的延迟差异。哪怕Bank冲突会让单次内存访问慢32倍，同步操作的固定大开销也会把这个差异稀释到几乎无法观测。

2. 计算量占比极低
   循环内的计算仅为shared_memory[index] += index * 10，属于极简单的算术操作，耗时可以忽略不计。整个kernel的运行时间几乎完全由同步和内存访问的固定成本决定，Bank冲突带来的额外延迟占比微乎其微。

3. 单线程初始化的瓶颈干扰
   共享内存由单个线程初始化，这部分的开销可能已经远大于后续循环中Bank冲突带来的差异，进一步模糊了两种场景的性能区别。

修正测试方案的建议 #

要准确观测Bank冲突的性能差异，需调整代码：

• 移除循环内的__syncthreads()，改用批量连续的内存访问操作，让内存访问延迟成为运行时间的主导因素。
• 增加内存访问的迭代次数，或者适当提升计算量占比，放大Bank冲突带来的延迟差异。
• 改用多warp或多block的测试场景，避免单warp下的硬件调度优化掩盖冲突影响。
• 并行初始化共享内存，比如让每个线程负责初始化一部分数据，消除单线程初始化的开销干扰。

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