将2D数组映射到CUDA中的RGB到GreyScale程序中的块网格

Question

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

1. 1mp,192 cuda core/MP。
2. 最大线程数/块=1024.
3. 最大常驻曲数/mp=64.
4. 线程/块的最大尺寸=(1024,1024,64).
5. 最大网格尺寸尺寸=(2147483647,65535,65535).

我的疑问是-

1. 如何确定块大小和网格大小?
2. 如果我将块大小从(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());
}

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

_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，它将显示一个错误，因为每个块的最大线程数是1024。所以，我可以使用最大1024*1024像素。如果一个图像有更多的像素，我就不能使用这个。现在我已经得到了第一段代码(最上面的代码)的输出，但我无法理解它背后的逻辑。

Answer 1

在具有计算能力3.2的CUDA设备的技术规范中，例如Tegra TK1，我们可以看到一些与您所描述的性能结果相关的限制因素。例如:

每个多处理器的最大线程数:2048

每个块的最大线程数:1024

每个多处理机的最大驻留块数:16

每个多处理器的最大驻留扭曲数:64

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

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

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