How the SobelFilter in the cuda sdk works SobelFilter

I’m new to cuda. Recently I’m working on the examples of the cuda sdk. The SobelFilter in the sdk took me a lot of time, however, I cann’t understand how it works. So if there anyone who can explain the sourcecode of SobelFilter to me? Thanks!

On a quick glance it looks pretty straightforward to me. Can you tell a bit more about what part you find difficult to understand? Textures, function pointers, or Sobel filters?

My sdk’s version is 3.2. The src folder includes SobelFilter.cpp, SobelFilter_kernels.cu and SobelFilter_kernels.h.

What puzzles me is the function “global void SobelShared”

"short u = 4*blockIdx.x*BlockWidth;

short v = blockIdx.y*blockDim.y + threadIdx.y;

short ib;"

At the top of SobelFilter_kernels.cu, it “#define BlockWidth 80 #define SharedPitch 384”

What is the BlockWidth & SharedPitch ?

As I read it, BlockWidth is the width of a block of data that is loaded into shared memory, while SharedPitch is the Pitch (difference of addresses of two points on adjacent lines with the same x coordinate) in shared memory.

It’s not exactly clear to my why the pitch gets rounded up to a multiple of 64 bytes though, I have to admit. This looks like a measure to prevent bank conflicts on compute 1.x devices where shared memory is organized in 16 banks each 4 bytes wide. At first glance I can’t find a place where that would actually matter though.

Thanks for your answer~

So far, I’m working on the examples in the book “CUDA By Example”, some definitions in the CUDA SDK examples are not familiar to me. Thanks to your answer, I will think it over.