__shared__ float Ms[KERNEL_SIZE * KERNEL_SIZE];

__shared__ float Ns[BLOCK_SIZE + 4][BLOCK_SIZE + 4];

if(threadIdx.x < KERNEL_SIZE * KERNEL_SIZE)

Ms[threadIdx.x] = M.elements[threadIdx.x];

int Ty = threadIdx.x / (BLOCK_SIZE + 4);

int Tx = threadIdx.x - (Ty * (BLOCK_SIZE + 4));

int Basex = blockIdx.x * BLOCK_SIZE;

int Basey = blockIdx.y * BLOCK_SIZE;

float result = 0.0f;

if(Ty < (BLOCK_SIZE / 2) + 2)

{

if(Basex + Tx - 2 >= 0 && Basex + Tx - 2 < N.width)

{

if(Basey + Ty - 2 >= 0 && Basey + Ty - 2 < N.height){

Ns[Ty][Tx] = N.elements[(Basey + Ty - 2)* N.pitch + Basex + Tx - 2];

} else {

Ns[Ty][Tx] = 0.0f;

}

Ty += (BLOCK_SIZE / 2) + 2;

if(Basey + Ty - 2 >= 0 && Basey + Ty - 2 < N.height){

Ns[Ty][Tx] = N.elements[(Basey + Ty - 2)* N.pitch + Basex + Tx - 2];

} else {

Ns[Ty][Tx] = 0.0f;

}

} else {

Ns[Ty][Tx] = 0.0f;

Ty += (BLOCK_SIZE / 2) + 2;

Ns[Ty][Tx] = 0.0f;

}

}

__syncthreads();

	embed