使用 CUDA 的快速图像模糊

• 下载源码 - 566.79 KB

引言

我采用 Mario Klingemann 开发的现有图像模糊算法，来演示如何使用 CUDA 进行图像模糊。请访问 http://incubator.quasimondo.com 了解有关堆栈模糊算法的更多详细信息。我认为堆栈模糊已经是外观良好的最快模糊算法。在此示例中，对 CUDA 进行了一些小的代码更改，并且我们看到了 CUDA 如何加速性能。

背景

模糊图像始终是一项耗时的任务。模糊质量和处理速度不能总是兼顾。 CUDA 可能会帮助程序员解决这个问题。此代码已在带有 NVIDIA GeForce G210M 的 Windows 7 上进行了测试。

传统的堆栈模糊

堆栈模糊需要先处理图像行，然后再处理图像列。有两个 while 循环来连续处理图像的行和列。耗时的部分是用于图像行和列的外部 while 循环。因此，它们是要通过 CUDA 修改的目标。

参数

• unsigned long* pImage [in/out]：32 位图像缓冲区
• unsigned w [in]：图像宽度
• unsigned h [in]：图像高度
• unsigned r [in]：模糊级别

void stack_blur_rgba32(unsigned long* pImage, unsigned w, unsigned h, unsigned r)
{
    // ...
    // Process image rows, this outer while-loop will be parallel computed by CUDA instead
    do
    {
        // Get input and output weights
        do
        {
            // ...
        }
        while(++i <= r);
        
        // Blur image rows
        do
        {
            // ...        
        }
        while(++x < w);      
    }
    while(++y < h);
    
    // Process image columns, this outer while-loop will be parallel 
    // computed by CUDA instead
    do
    {
        // Get input and output weights
        do
        {
            // ...
        }
        while(++i <= r);
        
        // Blur image columns
        do
        {
            // ...
        }
        while(++y < h);      
    }
    while(++x < w);  
    // ...
}

在此示例中，通过 CPU 测试，处理时间为 0.063553 (ms)。接下来我们将看到 CUDA 如何提高性能。

使用 CUDA 的堆栈模糊

最重要的部分是堆栈缓冲区，它需要为每一行和每一列都有独立的缓冲区。由于线程并行运行，因此堆栈缓冲区必须分离并由各个行和列使用。除了 CUDA 代码之外，其余代码几乎没有变化。

参数

• uchar4* pImage [in/out]：32 位图像缓冲区
• uchar4* stack_data_horiz_ptr [in]：行的堆栈缓冲区
• uchar4* stack_data_vert_ptr [in]：列的堆栈缓冲区
• unsigned w [in]：图像宽度
• unsigned h [in]：图像高度
• unsigned r [in]：模糊级别
• bool bMapped [in]：支持“主机内存映射到设备内存”的标志

void StackBlur_GPU(uchar4* pImage, uchar4* stack_data_horiz_ptr, 
	uchar4* stack_data_vert_ptr, unsigned w, unsigned h, unsigned r, bool bMapped)
{
    unsigned div = ((r + r) + 1);
    unsigned divLenHoriz = (sizeof(uchar4) * div * h);
    unsigned divLenVert = (sizeof(uchar4) * div * w);
    unsigned sizeImage = ((w * h) << 2);
    uchar4* stack_dev_horiz_ptr = NULL;
    uchar4* stack_dev_vert_ptr = NULL;
    uchar4* pImage_dev_ptr = NULL;
    unsigned mul_sum = *(stack_blur8_mul + r);
    unsigned shr_sum = *(stack_blur8_shr + r);

    if (false == bMapped)
    {
        cudaMalloc((void**)&stack_dev_horiz_ptr, divLenHoriz);
        cudaMalloc((void**)&stack_dev_vert_ptr, divLenVert);
        cudaMalloc((void**)&pImage_dev_ptr, sizeImage);

        cudaMemcpy(pImage_dev_ptr, pImage, sizeImage, cudaMemcpyHostToDevice);
    }
    else
    {
        cudaHostGetDevicePointer((void**)&stack_dev_horiz_ptr, 
					(void*)stack_data_horiz_ptr, 0);
        cudaHostGetDevicePointer((void**)&stack_dev_vert_ptr, 
					(void*)stack_data_vert_ptr, 0);
        cudaHostGetDevicePointer((void**)&pImage_dev_ptr, (void*)pImage, 0);
    }

    StackBlurHorizontal_Device<<<(unsigned)::ceil((float)(h + 1) / 
	(float)_THREADS), _THREADS>>>(pImage_dev_ptr, stack_dev_horiz_ptr, 
	mul_sum, shr_sum, w, h, r);
    StackBlurVertical_Device<<<(unsigned)::ceil((float)(w + 1) / 
	(float)_THREADS), _THREADS>>>(pImage_dev_ptr, stack_dev_vert_ptr, 
	mul_sum, shr_sum, w, h, r);

    if (false == bMapped)
    {
        cudaMemcpy(pImage, pImage_dev_ptr, sizeImage, cudaMemcpyDeviceToHost);

        cudaFree( stack_dev_horiz_ptr );
        stack_dev_horiz_ptr = NULL;

        cudaFree( stack_dev_vert_ptr );
        stack_dev_vert_ptr = NULL;

        cudaFree( pImage_dev_ptr );
        pImage_dev_ptr = NULL;        
    }
}

参数

• unsigned long* lpHostBuf [in/out]：32 位图像缓冲区
• unsigned w [in]：图像宽度
• unsigned h [in]：图像高度
• unsigned r [in]：模糊级别
• unsigned bMapped [in]：支持“主机内存映射到设备内存”的标志

void StackBlur_Device(unsigned long* lpHostBuf, unsigned w, 
			unsigned h, unsigned r, bool bMapped)
{
    if (NULL == lpHostBuf)
    {
        return;
    }
    else if ((r < 1) || (w < 1) || (h < 1))
    {
        return;
    }
    else if (r > 254)
    {
        r = 254;
    }

    uchar4* stack_data_horiz_ptr = NULL;
    uchar4* stack_data_vert_ptr = NULL;
    unsigned div = ((r + r) + 1);
    unsigned divLenHoriz = (sizeof(uchar4) * div * h);
    unsigned divLenVert = (sizeof(uchar4) * div * w);

    if (false == bMapped)
    {
        stack_data_horiz_ptr = (uchar4*)malloc( divLenHoriz );
        stack_data_vert_ptr = (uchar4*)malloc( divLenVert );
    }
    else
    {
        cudaHostAlloc((void**)&stack_data_horiz_ptr, divLenHoriz, cudaHostAllocMapped);
        cudaHostAlloc((void**)&stack_data_vert_ptr, divLenVert, cudaHostAllocMapped);
    }

    StackBlur_GPU((uchar4*)lpHostBuf, stack_data_horiz_ptr, 
			stack_data_vert_ptr, w, h, r, bMapped);
    DebugPrintf("StackBlur_GPU: %x\n", cudaGetLastError());

    if (false == bMapped)
    {
        free( stack_data_horiz_ptr );
        stack_data_horiz_ptr = NULL;

        free( stack_data_vert_ptr );
        stack_data_vert_ptr = NULL;
    }
    else
    {
        cudaFreeHost( stack_data_horiz_ptr );
        stack_data_horiz_ptr = NULL;

        cudaFreeHost( stack_data_vert_ptr );
        stack_data_vert_ptr = NULL;
    }
}

在此示例中，通过 GPU 测试，处理时间仅为 0.000150 (ms)。使用 CUDA 处理的时间比传统方式快 300 倍或更多。

此代码有效吗？！

1. 检查此链接并查看哪些 NVIDIA GPU 支持 CUDA。
2. 此代码将输出调试消息。从此链接下载 dbgview。
3. 对于某些 NVIDIA GPU，每个块的较高线程数可能不起作用。请将 _THREADS 更改为较小的值，然后重新编译此代码。

关注点

结果显然表明，使用 CUDA 进行并行计算非常棒。

历史

• 2009 年 9 月 8 日：首次发布
• 2009 年 9 月 10 日：将每个块的线程数的值更改为 256，以适应大多数 NVIDIA GPU