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

【问题标题】：CUDA - Optimize mean of matrix rows calculation using shared memoryCUDA - 使用共享内存优化矩阵行计算的平均值
【发布时间】：2015-02-17 22:34:55
【问题描述】：

我正在尝试优化我的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);

【问题讨论】：

你的共享内存内核在我看来是不连贯的。所以我不确定我是否可以以此为起点来做任何明智的事情。解决这个问题的一个明智的方法是对每行进行共享内存并行减少（SO上有很多这样的例子，或者你可以谷歌“cuda并行减少”，或者有一个CUDA示例代码），计算总和。一旦得到每行的总和，就可以计算每行的平均值。然后将这个平均值广播到处理该行的每个线程（通过共享内存），然后减去。您需要为每行分配一个线程块。
@RobertCrovella 我已经为每行分配了一个线程块； 1024 行 == 1024 个线程块，每个块有 256 个线程。我大步穿过 512 个元素的行，一次 256 个元素。所以，我无法想象我哪里出错了？
@AMostMajestuousCapybara sm_30, Quadro K6000
并行缩减不涉及以固定步幅跨行。我建议您学习如何编写经典的并行归约，而不是试图想象您哪里出错了。（仅作为损坏的一个示例，您似乎没有意识到 sumAllPixelsInRow 是唯一的，即 每个线程 单独的局部变量）。您没有并行减少，并且您的代码与一个几乎没有相似之处。 Here's an SO question 这当然是相关的。
Here 是关于 CUDA 并行缩减的规范教程。

标签： c++ cuda

【解决方案1】：

2 块不足以使 GPU 使用饱和。您正在采用更多块的正确方法，但是，您正在使用 Kepler，我想提出一个根本不使用共享内存的选项。

从一个块中的 32 个线程开始（稍后可以使用 2D 块进行更改）使用这 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 个线程，每个线程都有部分和。接下来，您需要将它们全部加在一起。通过使用 shuffle 减少，您可以在不使用共享内存的情况下执行此操作（因为我们在 warp 中）。有关该外观的详细信息：http://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/ 您想要的是 shuffle warp reduce，但您需要将其更改为使用完整的 32 个线程。

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

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

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

祝你好运。

【讨论】：