用 CUDA 在 GPU 上并行一个简单的算法

Question

我有一个在 GPU 上计算 Local Binary Patterns 的 CUDA 函数。基本上，LBP 是对图像像素的计算，其中任何给定像素 (i,j) 的值取决于它的 8 个邻居的强度。

到目前为止一切顺利，代码如下：

//The kernel
__global__ void LBP(unsigned char *in, unsigned char *out, const int w, const int h)
{
    const unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;

    //Don't do edges!
    if(
             i < w              //first row
        ||   i >= (w * (h - 1)) // last row
        || !(i % w)             // first column
        ||  (i % w + 1 == w)    // last column
    )
    {
        out[i] = 0;
        return;
    }

    unsigned char
        code = 0,
        center = in[i];

    code |= (in[i-w-1] > center) << 7;
    code |= (in[i-w  ] > center) << 6;
    code |= (in[i-w+1] > center) << 5;
    code |= (in[i  +1] > center) << 4;
    code |= (in[i+w+1] > center) << 3;
    code |= (in[i+w  ] > center) << 2;
    code |= (in[i+w-1] > center) << 1;
    code |= (in[i  -1] > center) << 0;

    out[i] = code;
}

// A proxi function
void DoLBP(unsigned char *in, unsigned char *out, const int w, const int h)
{
    const int
        sz = w * h * sizeof(unsigned char);
    unsigned char
        *in_gpu,
        *out_gpu;

    cudaMalloc((void**)&in_gpu,  sz);
    cudaMalloc((void**)&out_gpu, sz);

    cudaMemcpy(in_gpu,  in,  sz, cudaMemcpyHostToDevice);
    cudaMemcpy(out_gpu, out, sz, cudaMemcpyHostToDevice);

    dim3 threadsPerBlock(1024); //Max
    dim3 numBlocks(w*h/threadsPerBlock.x + 1);
    LBP<<<numBlocks,threadsPerBlock>>>(in_gpu, out_gpu, w, h);

    cudaMemcpy(out, out_gpu, sz, cudaMemcpyDeviceToHost);

    cudaFree(in_gpu);
    cudaFree(out_gpu);
}

//The caller
int main()
{
    printf("Starting\n");

    const int
        w = 4000,
        h = 2000;
    unsigned char
        in[w*h],
        out[w*h];

    // Fill [in] with some data 

    DoLBP(in, out, w, h);


    // Use [out] data

    return 0;
}

图像作为 *unsigned char*s (array = [[row 1] [row 2] [row 3] ... [row n]]) 的单维数组传递给 GPU（它们是从 OpenCV 的 Mat 中提取的）

问题

此代码适用于相对较小的图像，它返回填充了正确值的输出数组但是当图像大小增加时，输出数组全部归零！

我的怀疑是图像数据溢出了一些 GPU 缓冲区或类似的东西。

我也不清楚 numberOfBlocks 和 threadsPerBlock 部分是如何工作的！如果你们中的任何人能就此提供一些基本的见解，我们将不胜感激。

（我在 CUDA 中就像 1 天大，所以可能有太多方法可以改进这个 sn-p 代码！）

Answer 1

1. 我建议将proper cuda error checking 添加到您的代码中。我相信您的内核进行了越界访问并且失败了。
2. 使用cuda-memcheck 运行您的代码，因为它有助于确定内核失败的原因。

3. 这些是在堆栈上进行的相当大的分配：

   const int
     w = 4000,
     h = 2000;
   unsigned char
     in[w*h],
     out[w*h];
   

   每个大约 8MB。这可能是个问题；它可能取决于系统。通常最好通过动态分配进行大分配，例如malloc。在我的特定系统上，由于未正确分配这些大型堆栈变量，我遇到了段错误。

4. 您的内核缺少适当的“线程检查”。起初我以为你在这方面做得很好：

   if(
        i < w              //first row
     ||   i >= (w * (h - 1)) // last row
     || !(i % w)             // first column
     ||  (i % w + 1 == w)    // last column
   )
   

   但这是个问题：

   out[i] = 0;
   return;
   

   如果你注释掉out[i] = 0; 行，你会有更好的运气。或者，如果您不喜欢将其注释掉，您可以这样做：

   if (i < (w*h)) out[i] = 0;
   

   问题是您的网格启动参数必然会创建“额外线程”：

   dim3 threadsPerBlock(1024); //Max
   dim3 numBlocks(w*h/threadsPerBlock.x + 1);
   

   如果您进行了适当的线程检查（您几乎可以这样做......），那么这不是问题。但是你不能让那些额外的线程写入无效的位置。

要解释每个块的线程和块数，通过算术工作可能会很有用。 cuda 内核启动具有关联的 grid。网格只是与内核启动相关的所有线程。线程将被分成块。所以网格等于启动的块数乘以每个块的线程数。你的情况是多少？此行表示您要求每个块有 1024 个线程：

    dim3 threadsPerBlock(1024); //Max

您要启动的块数由以下公式给出：

    dim3 numBlocks(w*h/threadsPerBlock.x + 1);

算术是：

    (w=4000)*(h=2000)/1024 = 7812.5 = 7812   (note this is an *integer* divide)

然后我们添加 1。所以你正在启动 7813 个块。那是多少个线程？

    (7813 blocks)*(1024 threads per block) = 8000512 threads

但是你只需要（并且只想要）8000000 个线程 (= w * h) 所以你需要一个线程检查来防止额外的 512 个线程试图访问out[i]。但是您的线程检查在这方面被破坏了。

最后一点，对我来说，让这段代码运行得更快的最明显方法是通过共享内存利用相邻操作中的数据重用。但首先要让您的代码正常工作。