Lucas Kanade Implementation
Hi,

I am a rookie in CUDA world. My aim is to get a working program of Lucas-Kanade optical flow (naive) algorithm.
As the first step, I began with the process of loading image content onto device array. I have a few doubts in this:

1. Am I supposed to go the texture way? By this I mean to say is, loading the image using cutLoadPPM4ub() and then doing an initialization of texture for this image which copies the image content onto device memory in cudaArray and binds it to a texture.

2. If above step is fine, then why do I get a 0 when I try to read the image content using tex2D()?

So, the flow is like this..

[codebox]cutLoadPPM4ub(image_path, (unsigned char **) &h_img, &width, &height);

//other method
texture<float4, 2> tex_1, tex_2;
texture<float4, 2, cudaReadModeElementType> rgbaTex_1, rgbaTex_2;
int size = width * height * sizeof(unsigned int);
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
cudaMallocArray ( &d_array_1, &channelDesc, width, height );
cudaMemcpyToArray( d_array_1, 0, 0, pImage, size, cudaMemcpyHostToDevice);
tex_1.addressMode[0] = cudaAddressModeClamp;
tex_1.addressMode[1] = cudaAddressModeClamp;
tex_1.filterMode = cudaFilterModePoint;
tex_1.normalized = true;
cudaBindTextureToArray(tex_1, d_array_1, channelDesc);


...

In kernel
...

unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
float4 img1_val = tex2D(tex_1, float(x), float(y));

...
Doing a printf() for any of img1_val x, y or z component gives a 0 always.[/codebox]

P.S. I've tried filling the cudaArray with arbitrary values and it doesn't show up as 0s.


3. Last doubt is, in one step I need to calculate the derivative of Intensity of a pixel w.r.t time(x and y) as well. So, is this derivative the difference of the individual components of float4 values? If not, how to calculate the derivative of an image's intensity w.r.t a variable say x,y or t.
More precisely, what's the best way to get the intensity of a pixel? Is binding texture the right way?



Thanks!

--
Nitin
Hi,



I am a rookie in CUDA world. My aim is to get a working program of Lucas-Kanade optical flow (naive) algorithm.

As the first step, I began with the process of loading image content onto device array. I have a few doubts in this:



1. Am I supposed to go the texture way? By this I mean to say is, loading the image using cutLoadPPM4ub() and then doing an initialization of texture for this image which copies the image content onto device memory in cudaArray and binds it to a texture.



2. If above step is fine, then why do I get a 0 when I try to read the image content using tex2D()?



So, the flow is like this..



[codebox]cutLoadPPM4ub(image_path, (unsigned char **) &h_img, &width, &height);



//other method

texture<float4, 2> tex_1, tex_2;

texture<float4, 2, cudaReadModeElementType> rgbaTex_1, rgbaTex_2;

int size = width * height * sizeof(unsigned int);

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);

cudaMallocArray ( &d_array_1, &channelDesc, width, height );

cudaMemcpyToArray( d_array_1, 0, 0, pImage, size, cudaMemcpyHostToDevice);

tex_1.addressMode[0] = cudaAddressModeClamp;

tex_1.addressMode[1] = cudaAddressModeClamp;

tex_1.filterMode = cudaFilterModePoint;

tex_1.normalized = true;

cudaBindTextureToArray(tex_1, d_array_1, channelDesc);





...



In kernel

...



unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;

unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

float4 img1_val = tex2D(tex_1, float(x), float(y));



...

Doing a printf() for any of img1_val x, y or z component gives a 0 always.[/codebox]



P.S. I've tried filling the cudaArray with arbitrary values and it doesn't show up as 0s.





3. Last doubt is, in one step I need to calculate the derivative of Intensity of a pixel w.r.t time(x and y) as well. So, is this derivative the difference of the individual components of float4 values? If not, how to calculate the derivative of an image's intensity w.r.t a variable say x,y or t.

More precisely, what's the best way to get the intensity of a pixel? Is binding texture the right way?







Thanks!



--

Nitin

#1
Posted 03/29/2010 07:47 PM   
Sorry for this really late reply (3 months) but if you are still playing around or anyone else wants Optical Flow (Lucas & Kanade) or Optical Flow (Horn & Schunk) implementation on CUDA then it's freely available in the new release of [b][url="http://www.cuvilib.com"]CUVI Lib[/url][/b] which is a CUDA Accelerated Vision and Imaging library. Play around with it and give your feedback on the [b][url="http://www.cuvilib.com/forums"]forums[/url][/b]
Sorry for this really late reply (3 months) but if you are still playing around or anyone else wants Optical Flow (Lucas & Kanade) or Optical Flow (Horn & Schunk) implementation on CUDA then it's freely available in the new release of CUVI Lib which is a CUDA Accelerated Vision and Imaging library. Play around with it and give your feedback on the forums

Best,

Salman Ul Haq,

TunaCode Pvt. Limited

Developers of GPU based Computing solutions

(CUDA for Vision and Imaging Library)

W: www.cuvilib.com

E: salman@tunacode.com



GPU: NVIDIA Tesla C2050

CPU: Core2Duo 2.33GHz 2.5GB DDR2 RAM

CUDA Toolkit v4.0RC

#2
Posted 08/01/2010 02:35 PM   
Scroll To Top