在GPU上使用CUDA并行简单算法

Question

我有一个CUDA函数,可以计算GPU上的局部二进制模式.基本上,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天,因此可能有太多方法可以改进这段代码!)

Answer 1

1. 我建议在代码中添加适当的cuda错误检查.我相信你的内核正在进行越界访问和失败.
2. 运行代码cuda-memcheck,因为它将有助于确定内核失败的原因.

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

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

   每个大约8MB.那可能是个问题; 它可能与系统有关.通过动态分配进行大量分配通常会更好malloc.在我的特定系统上,由于没有正确分配这些大的堆栈变量,我得到了一个seg错误.

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内核启动具有关联的网格.网格只是与内核启动相关的所有线程.线程将被分成块.因此,网格等于每个块的线程启动的块数.在你的情况下有多少？这一行表示你要求每个块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].但是你的线程检查在这方面被打破了.

最后要说明的是,让代码运行得更快的最明显的方法是通过共享内存利用相邻操作中的数据重用.但首先要让你的代码正常工作.