启动更多的 CUDA 线程块是否比启动更少的线程块有更多的开销？答案

【问题标题】：Does launching more CUDA thread blocks have more overhead than launching fewer?启动更多的 CUDA 线程块是否比启动更少的线程块有更多的开销？
【发布时间】：2021-08-27 07:06:43
【问题描述】：

我正在用 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 块是一样的？

【问题讨论】：

试试/阅读这个：stackoverflow.com/questions/21332040/…（只需用总和替换演员表）。 for 循环的顺序会产生影响。

标签： c++ multithreading cuda gpu

【解决方案1】：

这个：

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

不是正确的grid-stride loop 内核设计。它可以称为块跨度循环。因此，当您使用此内核启动超过 1 个块时，每个块将精确地执行相同的操作。

我并不是说他们会“共同努力”完成任务；我的意思是每个块将单独完成任务。如果您启动 2 个区块，您将完成任务两次。

因此，试图从中推断出任何有关“开销”的信息都是不明智的。随着区块数量的增加您也在增加正在完成的工作量。

如果您进行了适当的网格步长循环内核设计，您会发现性能会提高到 GPU 饱和点，之后，随着网格大小（块数）的增加，您会发现性能几乎没有变化。

预计启动额外的区块会产生一些开销。断定某事物的成本为零是不合理的。然而，这种开销通常很小，这就是为什么一旦超过饱和点，grid-stride loop 的性能大致保持不变。饱和度的简单定义是足够的线程来为您正在运行的 GPU 中的每个 SM 提供完整的补充。

【讨论】：