Lucas Kanade Implementation


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.

  1. 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?



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