GPU线程占用率和同步停顿有什么关系？

Question

我正在编写一个带有内部循环的 CUDA 内核，大致如下所示：

for (int i = 0; i < NUM_ITERATIONS; i++)
{
  // read global memory, write shared memory
  __syncthreads();
  // read shared memory, do math
  __syncthreads();

}

为了性能，我想最小化线程等待其他线程到达屏障同步所花费的总时间。每个块的线程数是否会影响线程在屏障处等待的平均时间？所有线程等待的总时间？如果我的内核占用率较低（即每个线程有很多寄存器，每个块的线程数较低）怎么办？在这种情况下，是否有任何策略可以帮助减少同步等待？

Answer 1

每个块的线程数会影响线程在屏障处等待的平均时间吗？

有可能。每个块更多的线程肯定会增加最坏的情况。假设您在 RTX 3080 上每个块运行 512 个线程（CUDA 8.6，每个 SM 1536 个线程），那么三个扭曲（每个块一个）可能会因迟到而阻止同一 SM 上的其他 45 个扭曲前进。

你应该做你自己的基准测试，但上次我测试这个时，我最终使用了更小的块，即使它稍微降低了内存效率。

所有线程等待的总时间？如果我的内核占用率较低（即每个线程有很多寄存器，每个块的线程数较低）怎么办

低占用率可能会减少每个线程的等待时间，因为线程必须花费更少的时间来等待执行单元可用。但这通常是不可取的，因为您无法以这种方式充分利用计算性能。如果您希望每个线程的延迟较低，请留在 CPU 上。

在这种情况下有什么策略可以帮助减少同步等待吗？

您可以通过使用双缓冲来使用更多共享内存来消除这些障碍之一。而不是这样做：

__global__ void kernel()
{
  __shared__ float data[N];
  for (int i = 0; i < NUM_ITERATIONS; i++)
  {
    data[threadIdx.x] = load_global();
    __syncthreads();
    float loc = data[y];
    __syncthreads();
  }
}

这样做：

__global__ void kernel()
{
  __shared__ float data[2][N];
  for (int i = 0; i < NUM_ITERATIONS; i++)
  {
    data[i & 1][threadIdx.x] = load_global();
    __syncthreads();
    float loc = data[i & 1][y];
  }
}