CUDA 共享内存广播和 __syncthreads 行为

Question

我遇到了一个奇怪的问题，至少对我来说它看起来很奇怪，我希望有人能够阐明它。我有一个 CUDA 内核，它依赖于共享内存来进行快速本地访问。据我所知，如果半束中的所有线程都访问同一个共享内存库，那么该值将被广播到束中的线程。此外，从多个 warp 访问同一个 bank 不会导致 bank 冲突，它们只会被序列化。牢记这一点，我创建了一个小内核来测试它(在我的原始内核中遇到问题之后)。这是片段:

#define NUM_VALUES 16
#define NUM_LOOPS  1024

__global__ void shared_memory_test(float *output)
{
    // Create some shared memory
    __shared__ int dm_delays[NUM_VALUES];

    // Loop over NUM_LOOPS
    float accumulator = 0;
    for(unsigned c = 0; c < NUM_LOOPS; c++)
    {
        // Force shared memory update
        for(int d = threadIdx.x; d < NUM_VALUES; d++)
            dm_delays[d] = c * d;

        // __syncthreads();
        for(int d = 0; d < NUM_VALUES; d++)
            accumulator += dm_delays[d];
}

    // Store accumulated value to global memory
    for(unsigned d = 0; d < NUM_VALUES; d++)
        output[d] = accumulator;
}

我用 16 维的 block 运行了这个(半扭曲，不是非常有效，但它只是为了测试目的)。所有线程都应该寻址同一个共享内存库，所以不应该有冲突。然而，事实似乎恰恰相反。我在 Visual Studio 2010 上使用 Parallel Nsight 进行此测试。

对我来说更神秘的是，如果我取消注释外循环中的 __syncthreads 调用，那么银行冲突的数量会急剧增加。

只是一些数字给你一个想法(这是一个包含一个 block 的网格，有 16 个线程，所以一个半扭曲，NUM_VALUES = 16，NUM_LOOPS = 1024):

• 没有 __syncthreads:4 个库冲突
• 使用 __syncthreads:4,096 次银行冲突

我在 GTX 670 上运行它，设置为 compute_capability 3.0

提前致谢

更新:有人指出，如果没有 __syncthreads，编译器会优化外循环中的 NUM_LOOPS 读取，因为 dm_delays 的值永远不会改变。现在，在这两种情况下，我都遇到了 4,096 个存储区冲突，这仍然不能很好地适应共享内存的广播行为。

Answer 1

由于 dm_delays 的值没有改变，这可能是编译器在 __syncthreads 不存在的情况下优化了对共享内存的 1024 次读取的情况。有了 __syncthreads，它可能会假设该值可能会被另一个线程更改，因此它会一遍又一遍地读取该值。