为什么更改块和网格大小会对运行时产生如此大的影响？

Question

我正在研究一些将RGBA图像转换为灰度的cuda 教程.但我无法弄清楚为什么改变它blockSize并gridSize进行X33时间的改进.

__global__
void rgba_to_greyscale(const uchar4* const rgbaImage,
                       unsigned char* const greyImage,
                       int numRows, int numCols)
{
    int i = blockIdx.x*numCols + threadIdx.x;
    float channelSum = .299f * rgbaImage[i].x + .587f * rgbaImage[i].y + .114f * rgbaImage[i].z;
    greyImage[i]= channelSum;
}

void your_rgba_to_greyscale(const uchar4 * const h_rgbaImage, uchar4 * const d_rgbaImage,
                            unsigned char* const d_greyImage, size_t numRows, size_t numCols)
{
  const dim3 blockSize(numCols, 1, 1);
  const dim3 gridSize(numRows, 1 , 1);
  rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);

  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}

当我如上所述:

const dim3 blockSize(numCols, 1, 1);
const dim3 gridSize(numRows, 1 , 1);

我明白了 Your code executed in 0.030304 ms

当我设置:

 const dim3 blockSize(1, 1, 1);
 const dim3 gridSize(numRows, numCols , 1);

并更新线程函数以使用新索引:

int i = blockIdx.x*numCols + blockIdx.y;

我得到Your code executed in 0.995456 ms.

1. 我希望它是相反的,因为gpu可以在第二个网格分割上分别计算所有像素它是否与缓存一致性问题有关？为什么我得到这些结果？
2. 理论上这个问题的网格和块大小最好是什么？是否有可能在运行时计算它？

供参考:

numRows = 313 numCols =557 

---

技术性能:

#uname -a && /usr/bin/nvidia-settings -v
    Linux ip-10-16-23-92 3.2.0-39-virtual #62-Ubuntu SMP Thu Feb 28 00:48:27 UTC 2013 x86_64 x86_64 x86_64 GNU/Linux

    nvidia-settings:  version 304.54  (buildmeister@swio-display-x86-rhel47-11)

Answer 1

建议不要使用网格/块配置.第一个是不可扩展的,因为每个块的线程数对于GPU是有限的,因此它最终会因较大的图像大小而失败.第二个是糟糕的选择,因为每个块只有1个线程,不建议使用,因为GPU占用率非常低.您可以通过CUDA Toolkit附带的GPU占用计算器进行验证.推荐的块大小应该是GPU warp大小的倍数(16或32),具体取决于GPU.

在您的情况下,2D网格和块大小的一般和可扩展的方法将是这样的:

const dim3 blockSize(16, 16, 1);
const dim3 gridSize((numCols + blockSize.x - 1)/blockSize.x, (numRows + blockSize.y - 1)/blockSize.y , 1);

如果您保持在设备的限制范围内,您可以将块大小从16 x 16更改为您喜欢的任何大小.对于计算能力为1.0到1.3的设备,允许每块最多512个线程.对于计算能力2.0以后的设备,此限制是每块1024个线程.

现在,网格和块是2维的,内核中的索引将被修改如下:

int i = blockIdx.x * blockDim.x + threadIdx.x; //Column
int j = blockIdx.y * blockDim.y + threadIdx.y; //Row

int idx = j * numCols + i;

//Don't forget to perform bound checks
if(i>=numCols || j>=numRows) return;

float channelSum = .299f * rgbaImage[idx].x + .587f * rgbaImage[idx].y + .114f *     rgbaImage[idx].z;
greyImage[idx]= channelSum;