c++ - CUDA - 使用共享内存优化矩阵行计算的平均值

Question

我正在尝试优化 512w x 1024h 图像中每一行的平均值的计算，然后从计算它的行​​中减去平均值。我写了一段代码，它在 1.86 ms 内完成，但我想降低速度。这段代码工作正常，但不使用共享内存，而是使用 for 循环。我想除掉他们。

__global__ void subtractMean (const float *__restrict__ img, float *lineImg, int height, int width) {

  // height = 1024, width = 512

  int tidy = threadIdx.x + blockDim.x * blockIdx.x; 

  float sum = 0.0f; 
  float sumDiv = 0.0f; 

  if(tidy < height) { 

      for(int c = 0; c < width; c++) { 

          sum += img[tidy*width + c];
      }
      sumDiv = (sum/width)/2;

      //__syncthreads(); 

      for(int cc = 0; cc < width; cc++) { 

          lineImg[tidy*width + cc] = img[tidy*width + cc] - sumDiv;
      }

  }

  __syncthreads(); 

我调用上面的内核使用:

subtractMean <<< 2, 512 >>> (originalImage, rowMajorImage, actualImHeight, actualImWidth);

不过，我写的下面这段代码是使用共享内存来优化的。但是，它没有按预期工作。对问题可能是什么有任何想法吗？

__global__ void subtractMean (const float *__restrict__ img, float *lineImg, int height, int width) {

  extern __shared__ float perRow[];

  int idx = threadIdx.x;    // set idx along x
  int stride = width/2; 

  while(idx < width) { 
      perRow[idx] = 0; 
      idx += stride; 
  }

  __syncthreads(); 

  int tidx = threadIdx.x;   // set idx along x
  int tidy = blockIdx.x;    // set idx along y

  if(tidy < height) { 
      while(tidx < width) { 
          perRow[tidx] = img[tidy*width + tidx];
          tidx += stride; 
      }
  }

  __syncthreads(); 

  tidx = threadIdx.x;   // reset idx along x
  tidy = blockIdx.x;    // reset idx along y

  if(tidy < height) { 

      float sumAllPixelsInRow = 0.0f; 
      float sumDiv = 0.0f; 

      while(tidx < width) { 
          sumAllPixelsInRow += perRow[tidx];
          tidx += stride;
      }
      sumDiv = (sumAllPixelsInRow/width)/2;

      tidx = threadIdx.x;   // reset idx along x

      while(tidx < width) { 

          lineImg[tidy*width + tidx] = img[tidy*width + tidx] - sumDiv; 
          tidx += stride;
      }
  }

  __syncthreads();  
}

调用共享内存函数时使用:

subtractMean <<< 1024, 256, sizeof(float)*512 >>> (originalImage, rowMajorImage, actualImHeight, actualImWidth);

Answer 1

2 个 block 几乎不足以饱和 GPU 使用。您正朝着使用更多 block 的正确方法前进，但是，您正在使用开普勒，我想提出一个根本不使用共享内存的选项。

从一个 block 中的 32 个线程开始(稍后可以使用 2D block 更改) 对于这 32 个线程，您应该按照以下方式做一些事情:

int rowID = blockIdx.x;
int tid   = threadIdx.x;
int stride= blockDim.x;
int index = threadIdx.x;
float sum=0.0;
while(index<width){
    sum+=img[width*rowID+index];
    index+=blockDim.x;
}

此时您将拥有 32 个线程，每个线程中都有一个部分和。接下来您需要将它们全部加在一起。您可以在不使用共享内存的情况下(因为我们在 warp 中)通过使用 shuffle reduction 来做到这一点。有关详细信息，请查看此处:http://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/你想要的是 shuffle warp reduce，但你需要更改它以使用完整的 32 个线程。

现在每个 warp 中的线程 0 具有每一行的总和，您可以将其除以转换为 float 的宽度，并使用 shfl 使用 shfl(average, 0) 将其广播到 warp 的其余部分;。 http://docs.nvidia.com/cuda/cuda-c-programming-guide/#warp-description

找到平均值并隐式和显式同步扭曲(使用 shfl)后，您可以使用减法以类似的方法继续。

可能的进一步优化是在 block 中包含多个 warp 以提高占用率，并在宽度上手动展开循环以提高指令级并行性。

祝你好运。