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

Question

我的Cuda程序获得了显着的性能提升(平均),具体取决于块的大小和块数; 其中"线程"的总数保持不变.(我不确定线程​​是否是正确的术语......但我将在这里使用它;每个内核的总线程数是(块数)*(块大小)).我制作了一些图表来说明我的观点.

但首先让我先解释一下我的算法是什么,但我不确定它是多么相关,因为我认为这适用于所有GPGPU程序.但也许我错了.

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

__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):

事实上,这是我用这些图表制作的专辑,每张图表代表不同的问题大小:图形专辑.

我正在运行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个线程.基于前面的讨论,当我们最大化线程补码(最多2048)时,SM将具有隐藏延迟(因此通常提供最大平均性能)的最大能力.

对于512个线程的块大小,我们可以在SM上准确地匹配这些线程块中的4个,然后它将具有2048个线程的补充,可以从中选择工作和延迟隐藏.

但是当您将线程块大小更改为576,4*576> 2048时,我们就无法在每个SM上容纳4个线程块.这意味着,对于该内核配置,每个SM将运行3个线程块,即2048个中的1728个线程.从SM的角度来看,这实际上比以前允许2048个线程的情况更糟糕,因此它可能是性能从512减少到576个线程的指标(就像它从448增加到512一样,这涉及到瞬时占用率的类似变化).

由于上述原因,当我们改变每个块的线程时,看到性能图表(如您所示的那样)并不罕见.

具有粒度(量化)效果的其他占用限制器可导致性能图中类似的峰值行为.例如,你的问题中没有足够的信息来推测每线程寄存器的使用情况,但占用限制器可能是每个线程使用的寄存器.当您改变线程补码时,您会发现每个SM可能同样具有不断变化的块驻留,这可能导致不同的占用(上下)以及不同的性能.

为了进一步深入研究,我建议您花些时间了解各种分析器的占用率,每个线程的寄存器和性能分析功能.有很多关于这些主题的信息; 谷歌是你的朋友,并注意上面评论中链接的问题/答案,作为一个合理的起点.要充分研究占用率及其对性能的影响,需要比您在此处提供的信息更多的信息.它基本上需要MCVE以及确切的编译命令行,以及您运行的平台和CUDA版本.编译器的每线程寄存器使用受所有这些事情的影响,其中大部分都没有提供.