Translating four nested loops into a CUDA kernel

Question

I'm writing CUDA program that adds blur effect onto BMP files. I wrote working program that does this on CPU, now I'm trying to convert the code to CUDA. This is the function I want to work on CUDA:

void blur(bitmap_header* hp, unsigned char *data)
{
  int xx,yy,x,y, avgB, avgG, avgR, ile;
  int blurSize = 5;
    for(xx = 0; xx < hp->width; xx++)
    {
      for(yy = 0; yy < hp->height; yy++)
    {
        avgB = avgG = avgR = 0;
        ile = 0;

        for(x = xx; x < hp->width && x < xx + blurSize; x++)
        {


            for(y = yy; y < hp->height && y < yy + blurSize; y++)
            {
                avgB += data[x*3 + y*hp->width*3 + 0];
                avgG += data[x*3 + y*hp->width*3 + 1];
                avgR += data[x*3 + y*hp->width*3 + 2];
                ile++;
            }
        }

        avgB = avgB / ile;
        avgG = avgG / ile;
        avgR = avgR / ile;

        data[xx*3 + yy*hp->width*3 + 0] = avgB;
        data[xx*3 + yy*hp->width*3 + 1] = avgG;
        data[xx*3 + yy*hp->width*3 + 2] = avgR;
    }
}
}

How do I convert this function to work on CUDA device? Every tutorial covers only one for loop and uses

int i = threadIdx.x

My previous question about this program: Blur effect on bitmap using C

EDIT

FULL CODE with CUDA edits:

#include <stdio.h>
#include <stdlib.h>
#include <Windows.h>


#pragma pack(push,1)
/* Windows 3.x bitmap file header */
typedef struct {
char         filetype[2];   /* magic - always 'B' 'M' */
unsigned int filesize;
short        reserved1;
short        reserved2;
unsigned int dataoffset;    /* offset in bytes to actual bitmap data */
} file_header;

/* Windows 3.x bitmap full header, including file header */
typedef struct {
file_header  fileheader;
unsigned int headersize;
int          width;
int          height;
short        planes;
short        bitsperpixel;  /* we only support the value 24 here */
unsigned int compression;   /* we do not support compression */
unsigned int bitmapsize;
int          horizontalres;
int          verticalres;
unsigned int numcolors;
unsigned int importantcolors;
} bitmap_header;
#pragma pack(pop)

__global__ void blur(bitmap_header* hp, unsigned char *data)
{
int xx,yy,x,y, avgB, avgG, avgR, ile;
int blurSize = 5;

xx = blockIdy.y * blockDim.y + threadIdx.y;
yy = blockIdx.x * blockDim.x + threadIdx.x;

if(xx >= hp->width || yy >= hp->height)
    return;


avgB = avgG = avgR = 0;
ile = 0;

for(x = xx; x < hp->width && x < xx + blurSize; x++)
{


    for(y = yy; y < hp->height && y < yy + blurSize; y++)
    {
        avgB += data[x*3 + y*hp->width*3 + 0];
        avgG += data[x*3 + y*hp->width*3 + 1];
        avgR += data[x*3 + y*hp->width*3 + 2];
        ile++;
    }
}

avgB = avgB / ile;
avgG = avgG / ile;
avgR = avgR / ile;

data[xx*3 + yy*hp->width*3 + 0] = avgB;
data[xx*3 + yy*hp->width*3 + 1] = avgG;
data[xx*3 + yy*hp->width*3 + 2] = avgR;
}

int filter(char* input, char *output)
{
//variable dec:
FILE *fp,*out;
bitmap_header* hp;
bitmap_header* d_hp;
unsigned char *data;
unsigned char *d_data;

//Open input file:
fp = fopen(input, "r");
if(fp==NULL)
    return 1;

//Read the input file headers:
hp=(bitmap_header*)malloc(sizeof(bitmap_header));
cudaMalloc( &d_hp, (sizeof(bitmap_header));
if(hp==NULL)
    return 1;

fread(hp, sizeof(bitmap_header), 1, fp);
cudaMemcpy(d_hp, &hp, (sizeof(bitmap_header), cudaMemcpyHostToDevice);

//Read the data of the image:
data = (unsigned char*)malloc(sizeof(char)*hp->bitmapsize);
cudaMalloc( &d_data, (sizeof(char)*hp->bitmapsize));

fseek(fp,sizeof(char)*hp->fileheader.dataoffset,SEEK_SET);
fread(data,sizeof(char),hp->bitmapsize, fp);
cudaMemcpy(d_data, &data, (sizeof(char)*hp->bitmapsize), cudaMemcpyHostToDevice);


dim3 block(16,16);
dim3 grid ( (hp->height + 15)/16, (hp->width + 15)/16 );

blur<<<grid,block>>>(d_hp, d_data);
cudaMemcpy(data, d_data, (sizeof(char)*hp->bitmapsize), cudaMemcpyDeviceToHost);

//Open output file:
out = fopen(output, "wb");
if(out==NULL)
{
    fclose(fp);
    free(hp);
    free(data);
    cudaFree(d_data);
    cudaFree(d_hp);
    return 1;
}

fwrite(hp,sizeof(char),sizeof(bitmap_header),out);

fseek(out,sizeof(char)*hp->fileheader.dataoffset,SEEK_SET);
fwrite(data,sizeof(char),hp->bitmapsize,out);

fclose(fp);
fclose(out);
free(hp);
free(data);

cudaFree(d_data);
cudaFree(d_hp);
return 0;
}

int main(int argc, char* argv[])
{
int frames;
int frame = 1;
char path[100000];

system("rd /s/q temp");
system("mkdir temp");

system("cls");
printf("Zapis wszystkich klatek do folderu temp.\n");
system("ffmpeg.exe -i test.mp4 -r 29.970 -vcodec bmp temp/%d.bmp");

printf("Ile jest klatek w folderze temp?\n");
scanf("%d", &frames);

for(frame = 1;frame<=frames;frame++)
{
    sprintf(path,"temp\\%d.bmp",frame);
    printf("Nakladam filtr na ");
    printf(path);
    printf("\n");
    filter(path,path);
}

system("cls");
printf("Wszystkie klatki do filmu mp4.\n");
system("ffmpeg -r 29.970 -i temp/%d.bmp  -c:v libx264 -preset slow -crf 21 temp/out.mp4");

system("cls");
printf("Wyciecie dzwieku z filmu do mp3\n");
system("ffmpeg -i test.mp4 -vn -ar 44100 -ac 2 -ab 192 -f wav temp/sound.wav");

system("cls");
printf("Polaczenie mp3 z mp4.\n");
system("ffmpeg -i temp/sound.wav -i temp/out.mp4 final.mp4");

system("cls");
printf("Delete ");
system("rd /s temp");
system("pause");

return 0;
}

FULL ERROR LIST:

Warning    1    warning : The 'compute_10' and 'sm_10' architectures are deprecated, and may be removed in a future release. C:\Users\Karpińscy\documents\visual studio 2012\Projects\blur\blur\nvcc    blur
Error    2    error : identifier "blockIdy" is undefined C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur
Error    3    error : expected a ")" C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur
Warning    4    warning : expression has no effect C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur
Error    5    error : expected a ")" C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur
Error    6    error : too few arguments in function call C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur
Error    7    error MSB3721: The command ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\bin\nvcc.exe" -gencode=arch=compute_10,code=\"sm_10,compute_10\" --use-local-env --cl-version 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin"  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\include"  -G    -maxrregcount=0  --machine 32 --compile -cudart static -I  -g   -DWIN32 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd  " -o Debug\kernel.cu.obj "C:\Users\Karpińscy\documents\visual studio 2012\Projects\blur\blur\kernel.cu"" exited with code 2. C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\V110\BuildCustomizations\CUDA 6.0.targets    597    9    blur
    8    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    33    1    blur
    9    IntelliSense: expected a ';' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    33    12    blur
    10    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    56    4    blur
    11    IntelliSense: expected a ';' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    56    9    blur
    12    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    57    4    blur
    13    IntelliSense: expected a ';' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    57    7    blur
    14    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    58    3    blur
    15    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    62    2    blur
    16    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    63    2    blur
    17    IntelliSense: identifier "xx" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    65    7    blur
    18    IntelliSense: identifier "yy" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    65    14    blur
    19    IntelliSense: identifier "avgB" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    65    36    blur
    20    IntelliSense: identifier "xx" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    66    7    blur
    21    IntelliSense: identifier "yy" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    66    14    blur
    22    IntelliSense: identifier "xx" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    67    7    blur
    23    IntelliSense: identifier "yy" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    67    14    blur
    24    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    68    1    blur
    25    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    94    5    blur
    26    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    95    2    blur
    27    IntelliSense: identifier "d_data" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    95    15    blur
    28    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    95    21    blur
    29    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    97    5    blur
    30    IntelliSense: identifier "fp" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    97    11    blur
    31    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    97    13    blur
    32    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    98    5    blur
    33    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    98    15    blur
    34    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    99    2    blur
    35    IntelliSense: identifier "d_data" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    99    13    blur
    36    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    99    19    blur
    37    IntelliSense: identifier "dim3" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    102    2    blur
    38    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    102    15    blur
    39    IntelliSense: identifier "dim3" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    103    2    blur
    40    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    103    34    blur
    41    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    105    2    blur
    42    IntelliSense: expected a ';' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    105    6    blur
    43    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    106    2    blur
    44    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    106    17    blur
    45    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    109    2    blur
    46    IntelliSense: identifier "output" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    109    14    blur
    47    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    110    2    blur
    48    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    122    5    blur
    49    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    122    14    blur
    50    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    123    5    blur
    51    IntelliSense: expected a ')' c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    123    16    blur
    52    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    125    5    blur
    53    IntelliSense: identifier "fp" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    125    12    blur
    54    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    126    5    blur
    55    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    127    5    blur
    56    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    128    5    blur
    57    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    130    2    blur
    58    IntelliSense: identifier "d_data" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    130    11    blur
    59    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    131    2    blur
    60    IntelliSense: identifier "d_hp" is undefined c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    131    11    blur
    61    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    132    5    blur
    62    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    133    1    blur
    63    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    158    2    blur
    64    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    161    2    blur
    65    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    162    2    blur
    66    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    164    2    blur
    67    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    165    2    blur
    68    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    166    2    blur
    69    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    168    2    blur
    70    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    169    2    blur
    71    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    170    2    blur
    72    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    172    2    blur
    73    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    173    2    blur
    74    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    174    2    blur
    75    IntelliSense: this declaration has no storage class or type specifier    c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    175    2    blur
    76    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    177    5    blur
    77    IntelliSense: expected a declaration c:\Users\Karpińscy\Documents\Visual Studio 2012\Projects\blur\blur\kernel.cu    178    1    blur

a relatively straightforward conversion to CUDA would involve replacing the outer 2 for-loops with a 2D grid (using x and y thread variables), while retaining the inner 2 for-loops in the kernel. Each thread then becomes responsible for calculating one output point. For a 3D example, take a look [here](http://stackoverflow.com/questions/14920931/3d-cuda-kernel-indexing-for-image-filtering/14926201#14926201). Rather than asking for a tutorial or asking for someone to write the code for you, it's good practice to make an attempt yourself and then ask for help with specific questions/issues. — Robert Crovella, May 17 '14 at 14:33
@RobertCrovella I tried writing some CUDA code. Please check out edited question, and I'm getting errors. Do you know what might be the problem? — Bartosz Karpiński, May 17 '14 at 19:08
Post a complete code. One that I can cut, copy, paste, and compile. And regarding the error from VS, it would be good to see the full error output. If the only thing you're getting is what is posted in your question, your VS settings need to be modified to give more verbose output in the console window at the bottom. Then you will see the actual error that caused `nvcc` to exit with code 2. Also, use [proper cuda error checking](http://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api) on all cuda API calls and kernel calls. — Robert Crovella, May 17 '14 at 19:54
@RobertCrovella There are more errors, but I ignored them beacuse there were also on hello world program that worked on my setup. Full code + full erro list is here: http://pastebin.com/V7u8bmLr I extract frames from mp4 movie using FFMPEG, but it's not needed for compiling. — Bartosz Karpiński, May 17 '14 at 19:58
Post a complete code in the question which is as short as possible but demonstrates the problem. I'm not asking for your whole project. It requires effort on your part to create a short complete reproducer. — Robert Crovella, May 17 '14 at 20:14

score 5 · Accepted Answer · answered May 17 '14 at 20:47

There are quite a few syntax errors in your code. Really they don't have anything to do with CUDA. Perhaps you should improve your basic C coding skills and understanding how to interpret compiler errors.

This line:

cudaMalloc( &d_hp, (sizeof(bitmap_header));

is missing a close parenthesis. You can't see that, and you also can't figure out the compiler error that indicates that?

Error    3    error : expected a ")" C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur

In cuda, there is no built-in variable blockIdy. Perhaps you meant blockIdx.y?

Error    2    error : identifier "blockIdy" is undefined C:\Users\Karpi˝scy\documents\visual studio 2012\Projects\blur\blur\kernel.cu    blur

Another compile error on this line:

cudaMemcpy(d_hp, &hp, (sizeof(bitmap_header), cudaMemcpyHostToDevice);

It should be this:

cudaMemcpy(d_hp, &hp, sizeof(bitmap_header), cudaMemcpyHostToDevice);

If you still need help after fixing those compile errors, post a new question with a complete but simple code that doesn't depend on reading images from files. Just create a simple test case and use that to validate the kernel. Add the error checking I mentioned and run your code with cuda-memcheck. If you refuse to do those things, I can't help you.

score 1 · Answer 2 · edited May 23 '17 at 12:21

there's a method called prefix sum that requires to precompute a matrix of running sums along both axes. this can be done very efficiently. once you have this matrix, you can calculate sums/averages across arbitrary patches (submatrices) in O(1) by a linear combination of four values from the prefix sum matrix.

here's a thread about it: How to process a integer matrix to find the average for the elements of any sub-rectangle in O(1)?

the last step will be to iterate over the input image and set the values to that of local averages.

a different viewpoint to your bluring is the following: you convolve your image with a simple square filter (or kernel). there are multiple examples on SO and also on the web about implementing the 2D convolution in CUDA.

Translating four nested loops into a CUDA kernel

2 Answers2

Linked