GPGPU：块大小对程序性能的影响，为什么我的程序在非常特定的大小下运行得更快？

Question

我的 Cuda 程序获得显着的性能提升（平均而言）取决于块的大小和块的数量；其中“线程”的总数保持不变。 （我不确定线程​​是否是正确的术语......但我将在这里使用它；每个内核的线程总数是（块数）*（块大小））。我做了一些图表来说明我的观点。

但首先请允许我先解释一下我的算法是什么，但是我不确定它的相关性，因为我想这是适用于所有 GPGPU 程序的东西。但也许我我错了。

基本上我会遇到逻辑上被视为二维数组的大型数组，其中每个线程从数组中添加一个元素，并将该值的平方添加到另一个变量，然后最后将值写入另一个数组，在每次读取期间，所有线程都以某种方式移动。这是我的内核代码：

__global__ void MoveoutAndStackCuda(const float* __restrict__ prestackTraces, float* __restrict__ stackTracesOut,
  float* __restrict__ powerTracesOut, const int* __restrict__ sampleShift,
  const unsigned int samplesPerT, const unsigned int readIns,
  const unsigned int readWidth, const unsigned int defaultOffset) {

  unsigned int globalId = ((blockIdx.x * blockDim.x) + threadIdx.x); // Global ID of this thread, starting from 0 to total # of threads

  unsigned int jobNum = (globalId / readWidth); // Which array within the overall program this thread works on
  unsigned int readIndex = (globalId % readWidth) + defaultOffset; // Which sample within the array this thread works on

  globalId = (jobNum * samplesPerT) + readIndex;  // Incorperate default offset (since default offset will also be the offset of
                                                  // index we will be writing to), actual globalID only needed for above two variables.

  float stackF = 0.0;
  float powerF = 0.0;

  for (unsigned int x = 0; x < readIns; x++) {

    unsigned int indexRead = x + (jobNum * readIns);

    float value = prestackTraces[readIndex + (x * samplesPerT) + sampleShift[indexRead]];

    stackF += value;
    powerF += (value * value);
  }

  stackTracesOut[globalId] = stackF;
  powerTracesOut[globalId] = powerF;
}

现在是这篇文章的重点，调用这段代码时

  MoveoutAndStackCuda<<<threadGroups, threadsPerGroup>>>(*prestackTracesCudaPtr,
    *stackTracesOutCudaPtr, *powerTracesOutCudaPtr,
    *sampleShiftCudaPtr, samplesPerT, readIns,
    readWidth, defaultOffset);

我所做的只是在 >> 中使用不同的 threadGroups 和 threadsPerGroup，其中 threadGroups.x * threadsPerGroup.x 保持不变。 （如前所述，这是一个一维问题）。

我将块大小增加了 64，直到达到 1024。我预计不会有任何变化，因为我认为只要块大小大于 32，我相信这是内核中 ALU 的数量，它就会运行得一样快尽可能。看看我制作的这张图表：

对于这个特定大小，线程总数为 5000 * 5120，例如，如果块大小为 64，则有 ((5000 * 5120) / 64) 个块。 由于某种原因，在 896、768 和 512 块大小时性能显着提升。为什么？

我知道这看起来是随机的，但该图中的每个点都是 50 次测试的平均值！

这是另一个图表，这次是线程总数为 (8000 * 8192) 的时间。这次的提升是 768 和 960。

再举一个例子，这次是针对比其他两个问题更小的作业（总线程数为 2000 * 2048）：

事实上，这是我用这些图表制作的专辑，每张图表代表不同大小的问题：graph album。

我正在运行这个Quadro M5000，它有 2048 个 Cuda 核心。我相信每个 Cuda Core 都有 32 个 ALU，所以我假设在任何给定时间可能发生的计算总数是 (2048 * 32)？

那么是什么解释了这些神奇的数字呢？我认为它可能是线程总数除以 cuda 核心数，或除以 (2048 * 32)，但到目前为止，我发现与专辑中所有图表的任何内容都没有相关性。我可以做另一项测试来帮助缩小范围吗？我想找出运行这个程序的块大小以获得最佳结果。

我也没有包括它，但我也做了一个测试，块大小从 32 减少了 1，事情变得指数级地变慢。这对我来说很有意义，从那时起，我们每组的本地线程比给定多处理器中的 ALU 少。

Answer 1

基于此声明：

我将块大小增加了 64，直到达到 1024。我预计不会有任何变化，因为我认为只要块大小大于 32，我相信这是内核中 ALU 的数量，它就会运行得一样快尽可能。

我想说有一个关于 GPU 的重要概念，您可能不知道：GPU 是一种“隐藏延迟”的机器。它们主要通过向它们暴露大量可用（并行）工作来隐藏延迟。这可以粗略地概括为“很多线程”。对于 GPU 来说，一旦你有足够的线程来覆盖“核心”或执行单元的数量，这就足够了，这是一个完全错误的想法。 不是。

作为（初学者）GPU 程序员，您应该忽略 GPU 中的内核数量。您需要 很多 线程。在内核级别和每个 GPU SM 上。

一般来说，当您为每个 SM 提供更多线程时，GPU 在执行其他有用工作时隐藏延迟的能力就会增加。这解释了所有图表中的总体趋势，即斜率通常从左到右向下（即平均性能通常会提高，因为您向每个 SM 提供更多暴露的工作）。

但是，这并没有解决高峰和低谷。 GPU 存在大量可能影响性能的架构问题。我不会在这里提供完整的治疗。但让我们来看一个案例：

为什么第一张图中的性能增加到 512 个线程，然后突然下降到 576 个线程？

这很可能是占用效应。 GPU 中的 SM 最多有 2048 个线程。根据前面的讨论，当我们将线程补充最大化时，SM 将具有最大的隐藏延迟能力（因此通常会提供最大的平均性能），最高可达 2048。

对于 512 个线程的块大小，我们可以在一个 SM 上恰好放置 4 个这样的线程块，然后它将有 2048 个线程的补充，可供选择用于工作和延迟隐藏。

但是当你将线程块大小更改为 576 时，4*576 > 2048，所以我们不能再在每个 SM 上放置 4 个线程块。这意味着，对于该内核配置，每个 SM 将运行 3 个线程块，即 2048 个线程中的 1728 个线程。从 SM 的角度来看，这实际上更糟，比之前允许 2048 个线程的情况，因此它可能是性能从 512 到 576 个线程下降的一个指标（就像它增加了一样）从 448 到 512，这涉及到瞬时入住率的类似变化。

由于上述原因，当我们改变每个块的线程时，看到像您展示的那样的性能图表并不少见。

具有精细（量化）效果的其他占用限制因素可能会导致性能图表中出现类似的峰值行为。例如，您的问题中没有足够的信息来推测每个线程的寄存器使用情况，但占用的限制可能是每个线程使用的寄存器。当您改变线程补充时，您会发现同样地，每个 SM 驻留的块补充也可能发生变化，这可能导致占用率变化（上下），从而导致性能变化。

为了进一步深入研究，我建议您花一些时间了解各种分析器的占用率、每个线程的寄存器以及性能分析功能。已经有很多关于这些主题的信息； google 是您的朋友，请注意上面 cmets 中链接的question/answers，这是一个合理的起点。要全面研究入住率及其对性能的影响，需要比您在此处提供的信息更多的信息。它基本上需要MCVE 以及精确的编译命令行，以及您运行的平台和 CUDA 版本。编译器的 registers-per-thread 使用受到所有这些因素的影响，其中大部分是你没有提供的。