CUDA filtering video capture

Accelerated Computing CUDA CUDA Programming and Performance
1

I am working on a school project involving capturing video and filtering
I am using Windows XP, VS 2010, CUDA 5.0

I have filtered the video frame by frame and splitting each frame into four channels
with

_ global _ void separate_Channels( unsigned char* red, unsigned char* green, unsigned char* blue, uchar4* source, int iw, int ih) {
int x = (blockIdx.xblockDim.x) + threadIdx.x;
int y = (blockIdx.yblockDim.y) + threadIdx.y;

 int pixel = iw*y + x;

if (x < iw -1 && y < ih -1) {

  red[pixel]    = source[pixel].x;
  green[pixel]  = source[pixel].y;
  blue[pixel]   = source[pixel].z;

}

}

and then recombining with

_ global _ void combine_Channels ( uchar4* dest, unsigned char* red, unsigned char* green, unsigned char* blue, int iw, int ih ) { 
  int x = blockIdx.x*blockDim.x + threadIdx.x; 
  int y = blockIdx.y*blockDim.y + threadIdx.y;

  int pixel = x + iw*y;

  unsigned char r = red[pixel];
  unsigned char g = green[pixel];
  unsigned char b = blue[pixel];

  dest[pixel] = make_uchar4(r, g, b,255);

}

the filter Im using is a convolution filter

_ global _ void filter_kernel (int iw, int ih, unsigned char* source, unsigned char* >dest) { 
int x = (blockIdx.x*blockDim.x) + threadIdx.x; 
int y = (blockIdx.y*blockDim.y) + threadIdx.y;

//operate only if in the correct boundaries

if (x > 0 && x < iw -1 && y > 0 && y < ih -1) { int pixel = iw*y + x;

float kernel = (float)(0.7source[iw(y+1) + (x-1)] + 0.7source[iw(y+1)+ (x)]+0.7source[iw(y+1) + (x+1)]+
0.9source[iw(y) + (x-1)] + 0.9source[iw(y) + >(x)] + >0.9source[iw(y) + (x+1)] +
0.8source[iw(y-1) + (x-1)] + 0.8source[iw(y-1) + (x)] + 0.8source[iw(y-1) + (x+1)]);

  if (kernel > 255){
      dest[pixel] = 254;
  }
  else if (kernel < 0) {
      dest[pixel] = 1;
  }
  else {
      dest[pixel] = kernel;
  }

}

When I run without the filter, the channels recombine fine only problem is slanted dashed lines that appear When I run with the filter, the channels no long recombines fine and you can see noticeable shadows from other channels. The slanted dashed lines also appears

this is an output Im getting now https://fbcdn-sphotos-a-a.akamaihd.net/hphotos-ak-ash4/1001230_4847611442099_1569920245_n.jpg

I have tackling this problem for the past week and I can’t find any solution
It is getting quite frustrating as I need to get this done soon. Please Help

– Raph

2

Do You filter the channels separatly?

3

Yes
I capture the video as uchar4*
then split them into three channels R,G and B
filter them separately
then recombine

I have found a problem with the filter that might be causing the mult-overlap
for example if I use filters |0 0 0| or even |0 1 0| it would produce a perfectly normal image
--------------------------------|0 1 0|----------|0 1 0|
--------------------------------|0 0 0|----------|0 0 0|

but if i change the filter to
|0 0 0|
|1 1 0| you can see two images
|0 0 0|

|0 0 0|
|1 1 1| you can see three images
|0 0 0|

why is this so ?

4

You know that when filtering an image You need to devide the sum of the weighed pixel from the kernel by the sum of the weights. I can’t see it in Your code. I mean deviding, as Your code states, by

30.9 + 30.8 + 3*0.7 = 2.7 + 2.4 + 2.1 = 7.2

MK

5

Yes thanks for pointing that out!
but the multi-image overlap error still persists

6

Try casting each sample to ‘float’ before doing anything else. I can’t say this will solve the problem but doing so seems right. Secondly, the ‘kernel’ is ‘float’, while the output buffer You’re putting it into is ‘unsigned char’. Some data may be lost in such asignment. I recommend changing each sample to normalized float, meaning casting and deviding by 256.0f, like so:

float sample = ((float)source[iw*(y+1) + (x-1)]) / 256.0f;

After that do the summing and weighing, and before passing the outcome to destination do the reverse operation:

dest[pixel] = (unsigned char)(kernel * 256.0f);

You can also use multiply instead of devide - 1/256 is 0.00390625.

Cheers,
MK

7

In addition to the others posts, you are also missing the last column and the last row of pixels when splitting each frame into different channels in the separate_Channels kernel:

if (x < iw -1 && y < ih -1)

It should be

if (x < iw && y < ih)

In addition, you should do the same ‘best practice’ when recombining the channels in the combine_Channels kernel, although the image is a power of two.

Best,
Pablo.

8

hmm Thanks
will try those!!

also is it possible that the filter is suppose to do a shift
just on a higher resolution video this is very noticeable ?

9

Good linear filter should not be image resolution dependent. It should work the same and have similar outcomes for any resolution. Much depends on the weights You use whlie filtering, though.