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

【问题标题】：Why does changing the block and grid sizes have such great impact on runtime?为什么更改块和网格大小会对运行时产生如此大的影响？
【发布时间】：2013-06-15 00:51:02
【问题描述】：

我正在处理一些 cuda tutorial 将 RGBA 图片转换为灰度。但我不明白为什么更改 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。

我希望它是相反的，因为 gpu 可以计算所有第二个网格分裂上的像素分别是否与缓存一致性问题？为什么我会得到这些结果？
从理论上讲，这个问题的最佳网格和块大小是多少？是否可以在运行时计算？

仅供参考：

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)

【问题讨论】：

快速回答是 NVIDIA GPU 在称为 warp 的线程组中调度、获取、分派和执行指令。对于计算能力为 1.* - 3.* 的设备，warp 大小为 32。如果您启动仅包含 1 个线程的块，硬件会将块分为 1 个 warp，其中包括 1 个活动线程和 32 个非活动线程。数学管道和 LSU 的执行效率将是 1/32（你看是 32X）。

标签： c++ c cuda gpu gpgpu

【解决方案1】：

不推荐使用网格/块配置。第一个是不可扩展的，因为每个块的线程数对于 GPU 来说是有限的，因此它最终会因为更大的图像大小而失败。第二个是一个糟糕的选择，因为每个块只有 1 个线程，不建议这样做，因为 GPU 占用率会非常低。您可以通过 CUDA 工具包中包含的GPU Occupancy Calculator 进行验证。建议的块大小应该是 GPU 扭曲大小的倍数（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 个线程。

现在，网格和块是二维的，内核内部的索引将被修改如下：

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;

【讨论】：