更低的占用率 - 更好的性能

Question

以下问题一直困扰着我。

使用两个不同的设备运行相同的内核，一个具有 1.3 的计算能力，另一个具有 2.0 的计算能力，在 1.3 中每个块的线程数更多（高占用率），但在 2.0 中，我获得了更好的性能。 2.0 的性能峰值似乎是每个块 16 个线程，占用率为 17% 任何小于或大于该点的任何东西都具有最差的性能。

因为这很可能是内核本身的性质造成的。

__global__ void
kernel_CalculateRFCH (int xstart, int ystart, int xsize,
          int ysize, int imxsize, int imysize, int *test, int *dev_binIm, int *per_block_results)
{
  int x2, y2, bin, bin2;
  __shared__ int s_pixels[blockDim.x*blockDim.y];  //this wouldn't compile in reailty

  int tx = threadIdx.x;
  int ty = threadIdx.y;
  int tidy = threadIdx.y + blockIdx.y * blockDim.y;
  int tidx = threadIdx.x + blockIdx.x * blockDim.x;

  if (xstart + xsize > imxsize)
    xsize = imxsize - xstart;
  if (ystart + ysize > imysize)
    ysize = imysize - ystart;

  s_pixels[tx * blockDim.y + ty] = 0;

  if (tidy >= ystart && tidy < ysize + ystart && tidx >= xstart && tidx < xsize + xstart)
{
      bin = dev_binIm[tidx + tidy * imxsize];

      if (bin >= 0)
    {
      x2 = tidx;
      y2 = tidy;

         while (y2 < ystart + ysize)
          {
          if (x2 >= xstart + xsize || x2 - tidx > 10)
             {
                  x2 = xstart;
                  y2++;
                  if (tidx - x2 > 10)
                   x2 = tidx - 10;
                  if (y2 - tidy > 10)
                   {
                      y2 = ystart + ysize;
                      break;
                   }
                   if (y2 >= ystart + ysize)
                      break;
              }

          bin2 = dev_binIm[x2 + y2 * imxsize];

           if (bin2 >= 0)
              {
               test[(tidx + tidy * imxsize) * 221 + s_pixels[tx * blockDim.y + ty]] = bin + bin2 * 80;
               s_pixels[tx * blockDim.y + ty]++;
              }
          x2++;
        }           
     }          

  } 

  for (int offset = (blockDim.x * blockDim.y) / 2; offset > 0; offset >>= 1)
    {
     if ((tx * blockDim.y + ty) < offset)
       {
         s_pixels[tx * blockDim.y + ty] += s_pixels[tx * blockDim.y + ty + offset];
       }
      __syncthreads ();
     }

   if (tx * blockDim.y + ty == 0)
     {
        per_block_results[blockIdx.x * gridDim.y + blockIdx.y] = s_pixels[0];

     }

}

我使用二维线程。

ptxas 信息：为“sm_10”编译入口函数“_Z20kernel_CalculateRFCHiiiiiiPiS_” ptxas 信息：使用了 16 个寄存器，128 字节 smem，8 字节 cmem[1] .

每个设备的每个案例中都显示了 16 个寄存器。

任何关于为什么会发生这种情况的想法都会非常有启发性。

Answer 1

除了上面的一般性评论之外，您的内核是一个非常特殊的情况，因为大多数线程根本不做任何工作。为什么不直接将xstart 和ystart 添加到tidx 和tidy 并选择更小的网格？您在较小块大小上的更好性能可能只是感兴趣区域如何分割成块的人工制品。

这也解释了为什么您会发现计算能力 1.x 设备与 CC 2.0+ 设备之间存在巨大差异。从 CC 2.0 开始，Nvidia GPU 在处理运行时在块之间差异很大的内核方面已经变得更好。
在计算能力 1.x 上，只有在所有当前运行的块都完成后才会安排新的块浪潮，而从 CC 2.0 开始，只要任何旧块完成，就会启动新块。