Q:

Shared memory performing worse

I am performing sobel edge detection with a horizontal mask of size 3x3 and a vertical mask of size 3x3. 
I first implemented the algorithm without using shared memory and now I tried implementing it using shared memory. The performance when shared memory was not used is better than when it was. Can anybody explain please.
Code
__global__ void image(int * in, int *out, int width){
int Mx[3][3] = {{-1,0,1},{-2,0,2},{-1,0,1}};
int My[3][3] = {{1,2,1},{0,0,0},{-1,-2,-1}};
__shared__ int sIn[t][t];
int sumX = 0, sumY = 0;

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

int r = threadIdx.y;
int c = threadIdx.x;
sIn[r][c] = in[(row*width) + col];
__syncthreads();

if(row == 0 || row == width-1 || col == 0 || col == width-1){
    out[row*width + col] = 0;
}else{
    for(int i = -1; i < 2; i++){
        for(int j = -1; j < 2; j++){
            int pixel;
            if(r == 0 || r == (t-1) || c == 0 || c == (t-1))
                pixel = in[(row + i) * width + (col + j)];
            else
                pixel = sIn[r + i][c + j];
            sumX += pixel * Mx[i+1][j+1];
            sumY += pixel * My[i+1][j+1];
        }
    }
    int ans = abs(sumX) + abs(sumY);
    if(ans > 255) ans = 255;
    if(ans < 0) ans = 0;
    out[row*width + col] = ans;
}
}

t is the tile size.
Kernel call:
image<<<dim3(width/t, width/t, 1),dim3(t, t, 1)>>>(d_input, d_output, width);

A:

I used apron pixels thus increasing the tile width so as to reduce global memory accesses when the 3x3 mask's boundary pixels were checked

