对网格和块尺寸的混淆

Question

我正在尝试在 Udacity 课程的第 1 课结束时解决问题，但我不确定我是否只是打错字或实际代码是否错误。

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)
{
    size_t totalPixels = numRows * numCols;
    size_t gridRows = totalPixels / 32;
    size_t gridCols = totalPixels / 32;
    const dim3 blockSize(32,32,1);
    const dim3 gridSize(gridCols,gridRows,1);
    rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);
    cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}

另一种方法是：

void rgba_to_greyscale(const uchar4* const rgbaImage, unsigned char* const greyImage, int numRows, int numCols)
{   
    int x = (blockIdx.x * blockDim.x) + threadIdx.x;
    int y = (blockIdx.y * blockDim.y) + threadIdx.y;
    uchar4 rgba = rgbaImage[x * numCols + y];
    float channelSum = 0.299f * rgba.x + 0.587f * rgba.y + 0.114f * rgba.z;
    greyImage[x * numCols + y] = channelSum;
}

错误信息如下：

libdc1394 error: failed to initialize libdc1394
Cuda error at student_func.cu:76
unspecified launch failure cudaGetLastError()
we were unable to execute your code. Did you set the grid and/or block size correctly?

然后，它说代码已经编译，

Your code compiled!
error output: libdc1394 error: Failed to initialize libdc1394
Cuda error at student_func.cu:76
unspecified launch failure cudaGetLastError()

第 76 行是第一个代码块的最后一行，据我所知，我没有更改其中的任何内容。第76行如下，

rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);

我实际上找不到cudaGetLastError() 的声明。

我主要关心的是我对设置网格/块尺寸的理解 + 第一种方法在像素位置的一维数组与我的线程之间的映射方面是否正确。

编辑： 我想我误解了什么。 numRows 是垂直的像素数吗？ numCols是水平方向的像素吗？

我的块由 8 x 8 个线程组成，每个线程代表 1 个像素？如果是这样，我假设这就是为什么我在计算gridRows 时必须除以 4，因为图像不是正方形的？我假设我也可以制作一个 2:1 列的块：行？

编辑 2： 我只是尝试将我的块更改为 2:1 的比例，因此我可以将 numRows 和 numCol 除以相同的数字，但它现在在底部和侧面显示空白区域。为什么底部和侧面都有空白区域。我没有改变 by grid 或 block 的 y 维度。

Answer 1

每个块处理 32*32 像素，并且有 (totalPixels / 32) * (totalPixels / 32) 个块，所以你处理 totalPixels ^ 2 个像素 - 这似乎是错误的

第一个错了，这应该是正确的：

const dim3 blockSize(32,32,1);

size_t gridCols = (numCols + blockSize.x - 1) / blockSize.x;
size_t gridRows = (numRows + blockSize.y - 1) / blockSize.y;

这是 2d 的一个非常常见的模式 - 你可以记住它

在示例图像大小不是 2 的幂，并且您希望块覆盖所有图像（甚至更多）

所以下一个必须是正确的： gridCols * blockSize.x >= numCols gridRows * blockSize.y >= numRows

您选择块大小并根据它计算覆盖所有图像所需的块数量

之后，在内核中，您必须检查您是否没有“超出图像”，以防大小错误

另一个问题是在内核中，它必须是(y * numCols + x)，而不是相反

内核：

int x = (blockIdx.x * blockDim.x) + threadIdx.x;
int y = (blockIdx.y * blockDim.y) + threadIdx.y;

if(x < numCols && y < numRows)
{
    uchar4 rgba = rgbaImage[y * numCols + x];
    float channelSum = 0.299f * rgba.x + 0.587f * rgba.y + 0.114f * rgba.z;
    greyImage[y * numCols + x] = channelSum;
}

调用代码：

const dim3 blockSize(4,32,1); // may be any

size_t gridCols = (numCols + blockSize.x - 1) / blockSize.x;
size_t gridRows = (numRows + blockSize.y - 1) / blockSize.y;

const dim3 gridSize(gridCols,gridRows,1);
rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);
cudaDeviceSynchronize(); 
checkCudaErrors(cudaGetLastError());

该死的，我觉得我做的事情更难理解（