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

Question

我的 Cuda 程序获得了显着的性能提升(平均)，这取决于 block 的大小和 block 的数量； “线程”的总数保持不变。 (我不确定线程​​是否是正确的术语......但我将在这里使用它；对于每个内核，线程总数是(# of blocks)*(block size))。我制作了一些图表来说明我的观点。

但首先请允许我先解释一下我的算法，但是我不确定它的相关性如何，因为我想这是适用于所有 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 保持不变。 (如前所述，这是一个一维问题)。

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

对于此特定大小，线程总数为 5000 * 5120，因此例如，如果 block 大小为 64，则有 ((5000 * 5120)/64) 个 block 。 出于某种原因， block 大小为 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)，但到目前为止，我没有发现与我相册中所有图表的任何相关性。我可以做另一项测试来帮助缩小范围吗？我想找出运行此程序以获得最佳结果的 block 大小。

我也没有包括它，但我也做了一个测试，其中 block 大小从 32 减少了 1，并且速度呈指数级下降。这对我来说很有意义，因为在给定的多处理器中，我们每组的本地线程比 ALU 少。

Answer 1

基于这个声明:

I incremented the block size by 64 until I reached 1024. I expected no change, because I figured as long as block size is greater than 32, which I believe is the # of ALUs in a core, it would run as fast as possible.

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

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

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

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

Why does performance in the first graph increase up to 512 threads, then suddenly decrease at 576 threads?

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

对于 512 个线程的 block 大小，我们可以在 SM 上恰好容纳 4 个这样的线程 block ，然后它将有 2048 个线程的补充，可以从中选择工作和延迟隐藏。

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

由于上述原因，当我们改变每个 block 的线程数时，经常会看到像您所显示的那样的性能图表。

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

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