Sobel Filter ith shared memory

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi ,

I have a problem with this code(Sobel Filter):

__global__ void SobelFilter(unsigned char* g_DataIn, unsigned char* g_DataOut, int width, int height)

{

   __shared__ unsigned char sharedMem[BLOCK_HEIGHT * BLOCK_WIDTH];

   float s_SobelMatrix[9];

s_SobelMatrix[0] = -1;

    s_SobelMatrix[1] = 0;

    s_SobelMatrix[2] = 1;

s_SobelMatrix[3] = -2;

    s_SobelMatrix[4] = 0;

    s_SobelMatrix[5] = 2;

s_SobelMatrix[6] = -1;

    s_SobelMatrix[7] = 0;

    s_SobelMatrix[8] = 1;

// Computer the X and Y global coordinates

   int x = blockIdx.x * TILE_WIDTH + threadIdx.x ;//- FILTER_RADIUS;

   int y = blockIdx.y * TILE_HEIGHT + threadIdx.y ;//- FILTER_RADIUS;

// Get the Global index into the original image

   int index = y * (width) + x;

// Perform the first load of values into shared memory

   int sharedIndex = threadIdx.y * blockDim.y + threadIdx.x;

   sharedMem[sharedIndex] = g_DataIn[index];

   __syncthreads();

int i, j, rows, cols, startCol, endCol, startRow, endRow;

rows = height;

cols = width;

startCol = 1;

endCol = cols - 1;

startRow = 1;

endRow = rows - 1;

// Go through all inner pixel positions

for(i=startRow; i<endRow; i++) {

for(j=startCol; j<endCol; j++) {

// sum up the 9 values to calculate both the direction x and direction y

float sumX = 0, sumY=0;

for(int dy = -FILTER_RADIUS; dy <= FILTER_RADIUS; dy++) {

for(int dx = -FILTER_RADIUS; dx <= FILTER_RADIUS; dx++) {

float Pixel = (float)(sharedMem[i*width + j + (dy * width + dx)]); 	

sumX += Pixel * s_SobelMatrix[(dy + FILTER_RADIUS) * FILTER_DIAMETER +

(dx+FILTER_RADIUS)];

sumY += Pixel * s_SobelMatrix[(dx + FILTER_RADIUS) * FILTER_DIAMETER +

(dy+FILTER_RADIUS)];

}

}

g_DataOut[i*width + j] = (abs(sumX) + abs(sumY)) > EDGE_VALUE_THRESHOLD ? 255

: 0;

}

}

}

when I use shared memory,it doesn’t work(output image is a black image)

if I use DataIn (instead of shared memory), it generates the correct image

Can anyone tell me what’s wrong with shared memory here?

thanks

2

Any suggestions?

3

I notice you use BLOCK_HEIGHT and BLOCK_WIDTH in your shared memory declaration, while later TILE_WIDTH and TILE_HEIGHT are used for apparently the same purpose. Could this be related to the problem?

Do you check error codes to make sure the kernel is executed correctly?

4

It was solved.

I used “sharedIndex” instead of indexes for s_Sobel (“(dx + FILTER_RADIUS)” and “(dy + FILTER_RADIUS)”)

float Pixel = (float)(sharedMem[sharedIndex + (dy * blockDim.x + dx)]);

I also shouldn’t have wrote 2 “for” loops (variables “i” and “j”) and should have wrote it just to satisfy the limits

if((threadIdx.x >= FILTER_RADIUS) && (threadIdx.x < (BLOCK_WIDTH - FILTER_RADIUS))&&(threadIdx.y >= FILTER_RADIUS) && (threadIdx.y < (BLOCK_HEIGHT - FILTER_RADIUS)))
5

The SDK sample sobelFilter also implements Sobel.