c++ - 启动更多的 CUDA 线程 block 是否比启动更少的线程 block 有更多的开销？

Question

我正在用 CUDA 做一些实验，我注意到启动相同的基本内核:

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

与在不太大的阵列上仅启动一个块(总是有 512 个线程的块)相比，使用更多线程块有时会导致可执行文件的整体执行速度更慢。请注意，我一直在等待
在 CPU 上，这与创建线程的开销小于拥有更多线程所能产生的任何优势有关。
然而，在 GPU IIRC 上，我们没有通常意义上的线程，但我们只是拥有不同的物理内核，否则这些内核将不会被使用。我什至不认为这可能是内存问题，因为数据传输到 GPU 的时间不会改变，但也许我使用的统一内存正在做一些我不完全理解的事情。
所以我想知道:在 CUDA 中启动更多线程和线程块是否有更多开销？或者从 GPU 的角度启动 1 个块或 128 个块是否相同？

Answer 1

这个:

int index = threadIdx.x;
int stride = blockDim.x;
for (int i = index; i < n; i += stride)

不合适 grid-stride loop内核设计。它可以称为块步循环。因此，当您使用此内核启动 1 个以上的块时，每个块都将做完全相同的事情。
我不是说他们会“一起工作”来完成任务；我的意思是每个块将单独完成任务。如果您启动 2 个块，您将完成两次完成任务的工作。
因此，试图从中推断出任何有关“开销”的信息是不明智的。随着您增加块的数量，您也增加了正在完成的工作量。
如果您进行了适当的网格步长循环内核设计，您会发现性能提高到 GPU 饱和点，之后，随着网格大小(块数)的增加，您会发现性能几乎没有变化。
预计启动额外块会产生一些开销。断定某物的成本为零是不合理的。然而，这种开销通常非常小，这就是为什么一旦超过饱和点，网格步幅循环的性能就会大致保持不变。饱和度的简单定义是足够的线程为您运行的 GPU 中的每个 SM 提供完整的补充。