Click here to Skip to main content
Click here to Skip to main content

Fast Image Blurring with CUDA

By , 10 Sep 2009
Rate this:
Please Sign up or sign in to vote.

Introduction

I take an existing image blurring algorithm, developed by Mario Klingemann, to demonstrate the way to do image blurring with CUDA. Please visit http://incubator.quasimondo.com for more details about the stack blurring algorithm. I reckon that Stack Blur is already the fastest blur with a good looking algorithm. In this sample, there are some minor code changes with CUDA for this algorithm and we see how CUDA can speed up the performance.

Background

Blur image which is always a time consuming task. Blurring quality and processing speed cannot always have good performance for both. CUDA might help programmers resolve this issue. This code is tested on Windows 7 with NVIDIA GeForce G210M.

Stack Blur in a Conventional Way

Stack Blur needs to process image rows first and then columns. There are two while loops to process rows and columns of image consecutively. The time consuming parts are the outer while loops for image rows and columns. Therefore, they are the targets to be modified by CUDA.

Parameters

  • unsigned long* pImage [in/out]: 32-bit image buffer
  • unsigned w [in]: Image width
  • unsigned h [in]: Image height
  • unsigned r [in]: Blur level
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);  
    // ...
}

In this sample, processing time is 0.063553 (ms), tested by CPU. We will then see how performance can be sped up by CUDA.

Stack Blur with CUDA

The most significant parts are stack buffers which need to have independent buffers for each row and column. Because threads are running in parallel, stack buffers have to be separated and used by individual rows and columns. The rest of the code has nothing much changed except for the CUDA codes.

Parameters

  • uchar4* pImage [in/out]: 32-bit image buffer
  • uchar4* stack_data_horiz_ptr [in]: Stack buffer for rows
  • uchar4* stack_data_vert_ptr [in]: Stack buffer for columns
  • unsigned w [in]: Image width
  • unsigned h [in]: Image height
  • unsigned r [in]: Blur level
  • bool bMapped [in]: Flag of support of "host memory mapping to device memory"
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;        
    }
}

Parameters

  • unsigned long* lpHostBuf [in/out]: 32-bit image buffer
  • unsigned w [in]: Image width
  • unsigned h [in]: Image height
  • unsigned r [in]: Blur level
  • unsigned bMapped [in]: Flag of support of "host memory mapping to device memory"
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;
    }
}

In this sample, the processing time is only 0.000150 (ms), tested by GPU. The processing time with CUDA is 300x or more faster than the conventional way.

Does This Code Work?!

  1. Check this link and see what NVIDIA GPUs support CUDA.
  2. This code will spew debug messages. Download dbgview from this link
  3. Higher threads per block may not work for some NVIDIA GPUs. Please change _THREADS to a smaller value and re-compile this code.

Points of Interest

The result apparently tells that parallel computing with CUDA is amazing.

History

  • Sept. 8, 2009: Initial release
  • Sept. 10, 2009: Changed value of threads per block to 256, adapted to most of NVIDIA GPUs

License

This article, along with any associated source code and files, is licensed under The Code Project Open License (CPOL)

About the Author

ChaoJui
Software Developer (Senior) http://home.so-net.net.tw/lioucy
Taiwan Taiwan
I've been a coding guy for 15 years, using C/C++ and assembly. Also using database to do information presenation with graphics applications.

Comments and Discussions

 
BugThe timing on this is incorrect as noted above & VS2010 + CUDA 5.0 Support PinmemberJhomes4-Feb-13 14:36 
BugRe: The timing on this is incorrect as noted above & VS2010 + CUDA 5.0 Support [modified] Pinmemberido00015-Feb-13 0:06 
GeneralRe: The timing on this is incorrect as noted above & VS2010 + CUDA 5.0 Support PinmemberJhomes13-Mar-13 17:35 
Buggpu mode do not work for win7 x64 with latest driver\sdk\toolkit. Pinmembermrgloom22-Apr-12 8:02 
GeneralRe: gpu mode do not work for win7 x64 with latest driver\sdk\toolkit. Pinmemberh_nik00716-Nov-12 13:35 
GeneralMy vote of 5 Pinmembermbue14-Apr-11 11:44 
GeneralVS2010 PinmemberMajid Shahabfar9-Jul-10 9:56 
GeneralRe: VS2010 PinmemberJeffry Torres Alvarez2-Oct-13 1:53 
Questionif my laptop vdo card not nvdia,It work? Pinmemberpoi1191-May-10 9:35 
I will test your program for implement my project.Thank
GeneralElapsed time problems Pinmemberd3v14-Apr-10 8:58 
GeneralRe: Elapsed time problems [modified] PinmemberMember 770296525-Feb-11 9:46 
QuestionAny pixel differences between CUDA & non-CUDA image? PinmemberH. Gohel15-Sep-09 14:39 
AnswerRe: Any pixel differences between CUDA & non-CUDA image? PinmemberChaoJui16-Sep-09 2:06 
GeneralRe: Any pixel differences between CUDA & non-CUDA image? PinmemberJulekmen20-Oct-09 23:34 
GeneralRe: Any pixel differences between CUDA & non-CUDA image? PinmemberChaoJui22-Oct-09 4:39 
Generalcutil.h can't be found PinmemberCollin Jasnoch15-Sep-09 6:12 
GeneralRe: cutil.h can't be found [modified] PinmemberChaoJui15-Sep-09 23:49 
GeneralRe: cutil.h can't be found PinmemberCollin Jasnoch16-Sep-09 3:26 
GeneralRe: cutil.h can't be found PinmemberChaoJui16-Sep-09 4:14 
GeneralRe: cutil.h can't be found Pinmemberthorolfo7-Aug-11 0:08 
GeneralDoes not work with NVidia G8400 mobile (Win7, latest drivers) Pinmemberxliqz8-Sep-09 4:17 
GeneralRe: Does not work with NVidia G8400 mobile (Win7, latest drivers) [modified] PinmemberChaoJui8-Sep-09 4:35 
GeneralRe: Does not work with NVidia G8400 mobile (Win7, latest drivers) Pinmemberxliqz9-Sep-09 18:46 
GeneralRe: Does not work with NVidia G8400 mobile (Win7, latest drivers) Pinmemberxliqz10-Sep-09 5:34 
GeneralRe: Does not work with NVidia G8400 mobile (Win7, latest drivers) [modified] PinmemberChaoJui10-Sep-09 15:18 
GeneralRe: Does not work with NVidia G8400 mobile (Win7, latest drivers) Pinmemberxliqz12-Sep-09 22:03 
GeneralRe: Does not work with NVidia G8400 mobile (Win7, latest drivers) PinmemberChaoJui12-Sep-09 23:06 
GeneralNot work ! PinmemberApex Liu8-Sep-09 3:21 
GeneralRe: Not work ! [modified] PinmemberChaoJui8-Sep-09 3:48 

General General    News News    Suggestion Suggestion    Question Question    Bug Bug    Answer Answer    Joke Joke    Rant Rant    Admin Admin   

Use Ctrl+Left/Right to switch messages, Ctrl+Up/Down to switch threads, Ctrl+Shift+Left/Right to switch pages.

| Advertise | Privacy | Mobile
Web04 | 2.8.140415.2 | Last Updated 10 Sep 2009
Article Copyright 2009 by ChaoJui
Everything else Copyright © CodeProject, 1999-2014
Terms of Use
Layout: fixed | fluid