__global__

void SobelGPUShared(float *image, int height, int width, float *x_matrix,

float *y_matrix, float *output) {

extern __shared__ float arr[];

int i = blockIdx.x*blockDim.x+threadIdx.x;

int j = blockIdx.y*blockDim.y+threadIdx.y;

//arr[i*j] = image[i*j];

if(i < height && j < width)

{

float x_grad=0;

x_grad += x_matrix[0 * MATRIX_DIM + 0] *

GetValidPixelValue(image, height, width, i - 1, j - 1);

x_grad += x_matrix[0 * MATRIX_DIM + 2] *

GetValidPixelValue(image, height, width, i - 1, j + 1);

x_grad += x_matrix[1 * MATRIX_DIM + 0] *

GetValidPixelValue(image, height, width, i, j - 1);

x_grad += x_matrix[1 * MATRIX_DIM + 2] *

GetValidPixelValue(image, height, width, i, j + 1);

x_grad += x_matrix[2 * MATRIX_DIM + 0] *

GetValidPixelValue(image, height, width, i + 1, j - 1);

x_grad += x_matrix[2 * MATRIX_DIM + 2] *

GetValidPixelValue(image, height, width, i + 1, j + 1);

__syncthreads();

float y_grad = 0;

y_grad += y_matrix[0 * MATRIX_DIM + 0] *

GetValidPixelValue(image, height, width, i - 1, j - 1);

y_grad += y_matrix[0 * MATRIX_DIM + 1] *

GetValidPixelValue(image, height, width, i - 1, j);

y_grad += y_matrix[0 * MATRIX_DIM + 2] *

GetValidPixelValue(image, height, width, i - 1, j + 1);

y_grad += y_matrix[2 * MATRIX_DIM + 0] *

GetValidPixelValue(image, height, width, i + 1, j - 1);

y_grad += y_matrix[2 * MATRIX_DIM + 1] *

GetValidPixelValue(image, height, width, i + 1, j);

y_grad += y_matrix[2 * MATRIX_DIM + 2] *

GetValidPixelValue(image, height, width, i + 1, j + 1);

__syncthreads();

float magnitude =

sqrt(x_grad / 8 * x_grad / 8 +

y_grad / 8 * y_grad / 8); // normalize gradients by dividing by 8

if (magnitude > 1) { // clamp to 1

output[i * width + j] = 1;

} else if (magnitude > THRESHOLD) {

output[i * width + j] = magnitude;

} else {

output[i * width + j] = 0;

}

// Implement this function, define a GPU kernel above it

}

}

Given this Cuda C++ function, how can you divide the image array into tiles that will fit in shared memory? You only have to worry about the parallelization part and not the actual sobel filter calculations that are being done.