c++ - 在 RGB 到 GrayScale 程序中将二维数组映射到 CUDA 中的 block 网格

Question

我是 Cuda 编程新手。我正在尝试 RGB 到灰度转换。但是，我不知道如何选择 block 大小和网格大小。我遇到了这段代码并且它正确执行了。但我不明白 gridSize 是如何选择的。我正在使用 Tegra TK1 Gpu，它有-

1. 1 MP，192 cuda 核心/MP。
2. 线程/ block 的最大数量=1024。
3. 常驻经线的最大数量/mp=64。
4. 线程/ block 的最大维度大小=(1024,1024,64)。
5. 网格大小的最大维度=(2147483647,65535,65535)。

我的疑惑是-

1. 如何确定 block 大小和网格大小？
2. 如果我将 block 大小从 (16,16,1) 更改为 (32,32,1)，则花费的时间会更长。这是为什么？

您还可以提供与此相关的任何优秀论文/书籍的链接吗？提前谢谢你。

这是代码-

_global__
void rgba_to_greyscale(const uchar4* const rgbaImage,
                       unsigned char* const greyImage,
                       int numRows, int numCols)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x; //Column
    int j = blockIdx.y * blockDim.y + threadIdx.y; //Row

    int idx = j * numCols + i;

    if(i>=numCols || j>=numRows) return;

    float channelSum = .299f * rgbaImage[idx].x + .587f * rgbaImage[idx].y + .114f *     rgbaImage[idx].z;
    greyImage[idx]= 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(16, 16, 1);
const dim3 gridSize((numCols + (blockSize.x-1)) /blockSize.x , (numRows +(blockSize.y-1)) /blockSize.y, 1);
rgba_to_greyscale<<<gridSize,blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);
cudaDeviceSynchronize(); 
checkCudaErrors(cudaGetLastError());
}

编辑- 我在使用上述代码之前使用的代码，将二维数组映射到 CUDA 中的 block 网格是-

_global__
void rgba_to_greyscale(const uchar4* const rgbaImage,
                       unsigned char* const greyImage,
                       int numRows, int numCols)
{
    int col = threadIdx.x;
    int row = blockIdx.x;
    int idx = col+row*numCols;
    int R = rgbaImage[idx].x;
    int G = rgbaImage[idx].y;
    int B = rgbaImage[idx].z;
    greyImage[idx] = 0.299f*R + 0.587f*G + 0.114f*B;
}

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());

}

我理解这段代码中的错误。这里的错误是，如果 numRows 和 numCols 大于 1024，它会显示错误，因为每个 block 的最大线程数是 1024。所以，我可以使用最大 1024*1024 像素。如果图像有更多像素，我就不能使用它。现在我得到了第一个代码(最上面的代码)的输出，但我无法理解其背后的逻辑。

Answer 1

在technical specification for CUDA devices使用 3.2 的计算能力，例如 Tegra TK1，我们可以看到一些与您描述的性能结果相关的限制因素。参见示例:

Maximum number of threads per multiprocessor: 2048

Maximum number of threads per block: 1024

Maximum number of resident blocks per multiprocessor: 16

Maximum number of resident warps per multiprocessor: 64

如果我们(我)可以假设除了最大线程数之外没有任何限制因素(内核不使用共享内存，我认为每个线程的寄存器数将少于 63 个)。

然后，对于 16 x 16 线程 block ，即 256 线程或 8 线程束，我们最多有 8 每个 SM 并发 block (受每个 SM 并发 warp 的最大数量限制)。如果将 block 的大小更改为 32 x 32(1024 线程或 32 线程)，并发 block 的最大数量将为 2。这可能是主要原因，因为第二种配置的执行时间更长。

block 大小的最佳配置通常有点棘手，它是基于反复试验的。默认情况下，我们(我)总是开始最大化占用率，然后尝试其他配置。