关于cuda块索引维度的问题答案

【问题标题】：Question on the dimension of cuda block indexing关于cuda块索引维度的问题
【发布时间】：2022-01-17 20:22:28
【问题描述】：

我认为以下 cuda 代码取自“使用 GPU 计算加速 MATLAB：带示例的入门”一书

int row = blockIdx.x * blockDim.x + threadIdx.x;
if (row < 1 || row > numRows - 1)
    return;

int col = blockIdx.y * blockDim.y + threadIdx.y;
if (col < 1 || col > numCols - 1)
    return;

应该是

int row = blockIdx.x * blockDim.x + threadIdx.x;
if (row < 0 || row > numRows - 1)
    return;

int col = blockIdx.y * blockDim.y + threadIdx.y;
if (col < 0 || col > numCols - 1)
    return;

我说的对吗？以下是使用从 MATLAB 调用的 cuda 代码进行图像卷积的完整代码。

#include "conv2Mex.h"

__global__ void conv2MexCuda(float* src,
                             float* dst,
                             int numRows,
                             int numCols,
                             float* mask)
{
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < 1 || row > numRows - 1)
        return;

    int col = blockIdx.y * blockDim.y + threadIdx.y;
    if (col < 1 || col > numCols - 1)
        return;

    int dstIndex = col * numRows + row;
    dst[dstIndex] = 0;
    int mskIndex = 3 * 3 - 1;
    for (int kc = -1; kc < 2; kc++)
    {
        int srcIndex = (col + kc) * numRows + row;
        for (int kr = -1; kr < 2; kr++)
        {
            dst[dstIndex] += mask[mskIndex--] * src[srcIndex + kr];
        }
    }
}

void conv2Mex(float* src, float* dst, int numRows, int numCols, float* msk)
{
    ...
    conv2MexCuda<<<gridSize, blockSize>>>...
    ...
}

【问题讨论】：

row 和 col 都不会小于 0。要找出为什么这个内核需要测试 < 1，我建议在纸上画一个小例子。原因不是 CUDA 特定的

标签： matlab cuda gpu mex

【解决方案1】：

我说的对吗？

我不认为你是对的。

内核代码中row 和col 索引的构造使得它们（跨网格中的线程）从0 到numRows-1 和从0 到numCols-1（可能更大，取决于关于实际网格大小，您没有显示）。

根据您显示的代码，掩码显然是一个 3x3 掩码，这意味着它充当当前 (row, col) 位置的模板，并扩展正负一行，正负一柱子。让我们仔细看看这里的索引对于 (row, col) = (0,0); 的情况。这是您根据您提议的变更允许执行的职位之一：

for (int kc = -1; kc < 2; kc++)
{
    int srcIndex = (col + kc) * numRows + row;
    for (int kr = -1; kr < 2; kr++)
    {
        dst[dstIndex] += mask[mskIndex--] * src[srcIndex + kr];

在外部 for 循环的第一次迭代中，kc 将为 -1，因此 srcIndex 为 (0-1)*numRows+0。假设 numRows 相当大，比如 256。所以 srcIndex 是 -1*256 或 -256。在内部 for 循环的第一次迭代中，kr 为 -1，因此访问 src 的计算索引为 -256-1 = -257。这几乎是不明智的。

如果有的话，上限对我来说看起来不正确。如果我们假设有效的图像索引范围是 0..numRows-1 和 0..numCols-1，那么我认为限制应该如下：

int row = blockIdx.x * blockDim.x + threadIdx.x;
if (row < 1 || row > numRows - 2)
    return;

int col = blockIdx.y * blockDim.y + threadIdx.y;
if (col < 1 || col > numCols - 2)
    return;

这似乎是classic computer science off-by-1 error。

【讨论】：