3

我遇到了一个奇怪的问题，至少对我来说它看起来很奇怪，我希望有人能够对此有所了解。我有一个 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 的块尺寸运行它（半个扭曲，效率不是很高，但仅用于测试目的）。所有线程都应该寻址同一个共享内存库，所以应该没有冲突。然而，事实似乎恰恰相反。我在 Visual Studio 2010 上使用 Parallel Nsight 进行此测试。

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

只是一些数字给你一个想法（这是一个包含一个有 16 个线程的块的网格，所以一个半扭曲，NUM_VALUES = 16，NUM_LOOPS = 1024）：

  • 没有__syncthreads：4个银行冲突
  • with __syncthreads : 4,096 银行冲突

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

先感谢您

更新：有人指出，__syncthreads由于 dm_delays 的值永远不会改变，因此编译器正在优化外部循环中没有 NUM_LOOPS 读取。现在，在这两种情况下，我都会遇到 4,096 次银行冲突，这仍然不能很好地与共享内存的广播行为配合使用。

4

1 回答 1

0

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

于 2013-03-25T15:35:51.180 回答