Question: #include #include #include #include #include #include #include #define MATRIX_DIM 3 #define THRESHOLD 0.1 void LoadImage(const char *filename, float **data, int *height, int *width); int SaveImage(const

#include

#include

#include

#include

#include

#include

#include

#define MATRIX_DIM 3

#define THRESHOLD 0.1

void LoadImage(const char *filename, float **data, int *height, int *width);

int SaveImage(const char *filename, float *image, int height, int width);

__host__ __device__

float GetValidPixelValue(float *image, int height, int width, int r, int c) {

// Prevent trying to read pixel locations outside the image

if (r < 0 || r >= height || c < 0 || c >= width)

return 0;

else

return image[r * width + c];

}

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

float *y_matrix, float *output) {

for (int i = 0; i < height; i++) {

for (int j = 0; j < width; j++) {

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);

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);

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;

}

}

}

}

__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

}

}

__global__

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

float *y_matrix, float *output) {

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

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

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);

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);

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

}

}

double WallTime() { // returns time in MS as a double

struct timeval tv;

gettimeofday(&tv, 0);

return tv.tv_sec * 1000.0 + tv.tv_usec / 1000.0;

}

void CheckResults(int height, int width, float *gpu_result, float *cpu_result){

for (int i = 0; i < height * width; i++) {

if (fabs(gpu_result[i] - cpu_result[i]) > 0.01) {

printf(" Mismatch at index %d: GPU: %f CPU %f ", i, gpu_result[i], cpu_result[i]);

return;

}

}

printf(" Ok!!! ");

}

int main(int argc, char *argv[]) {

float *data;

double start, end;

int height, width;

LoadImage("input.bmp", &data, &height, &width);

// CPU/GPU Data

float *x_matrix;

float *y_matrix;

float *image;

float *cpu_result;

float *gpu_result;

cudaMallocManaged(&image, sizeof(float) * height * width);

cudaMallocManaged(&x_matrix, sizeof(float) * MATRIX_DIM * MATRIX_DIM);

cudaMallocManaged(&y_matrix, sizeof(float) * MATRIX_DIM * MATRIX_DIM);

cudaMallocManaged(&cpu_result, sizeof(float) * height * width);

cudaMallocManaged(&gpu_result, sizeof(float) * height * width);

memcpy(image, data, sizeof(float) * height * width);

// Sobel x kernel matrix (already flipped)

x_matrix[0] = -1;

x_matrix[1] = 0;

x_matrix[2] = 1;

x_matrix[3] = -2;

x_matrix[4] = 0;

x_matrix[5] = 2;

x_matrix[6] = -1;

x_matrix[7] = 0;

x_matrix[7] = 1;

// Sobel y kernel matrix (already flipped)

y_matrix[0] = -1;

y_matrix[1] = -2;

y_matrix[2] = -1;

y_matrix[3] = 0;

y_matrix[4] = 0;

y_matrix[5] = 0;

y_matrix[6] = 1;

y_matrix[7] = 2;

y_matrix[8] = 1;

start = WallTime();

SobelCPU(image, height, width, x_matrix, y_matrix, cpu_result);

end = WallTime();

printf("CPU Runtime = %f msecs ", end - start);

dim3 threadsPerBlock(16,16);

dim3 numBlocks(height/threadsPerBlock.x, width/threadsPerBlock.y);

cudaMemcpy(image, data, sizeof(float)*height*width,cudaMemcpyHostToDevice);

SobelGPUShared<<>>(image, height, width, x_matrix, y_matrix, gpu_result);

cudaMemcpy(gpu_result,gpu_result, sizeof(float)*height*width,cudaMemcpyDeviceToHost);

cudaDeviceSynchronize();

CheckResults(height, width, gpu_result, cpu_result);

SaveImage("cpu_result.bmp", cpu_result, height, width);

SaveImage("gpu_result.bmp", gpu_result, height, width);

free(data);

cudaFree(image);

cudaFree(cpu_result);

cudaFree(gpu_result);

cudaFree(x_matrix);

cudaFree(y_matrix);

return 0;

}

Given this piece of Cuda code, how would you fix the shared memory sobel filter function SobelGPUShared so that the image is divided up into tiles that fit into shared memory?

Step by Step Solution

There are 3 Steps involved in it

1 Expert Approved Answer
Step: 1 Unlock

To fix the SobelGPUShared function so that the image is divided into tiles that fit into shared memory the key is to load a tile a small block of the image into shared memory perform the Sobel filteri... View full answer

blur-text-image
Question Has Been Solved by an Expert!

Get step-by-step solutions from verified subject matter experts

Step: 2 Unlock
Step: 3 Unlock

Students Have Also Explored These Related Databases Questions!

Q:

C. Progamming. Implement Actions and Functions I have to implement the function and action: tError bookTable_sortedAdd(tBookTable *tabBook, tBook book) so given a table of books sorted according the...

Q:

Hello all. So I'm new to C++ and I've spent 3 days trying to make the code below compile!! If anyone can help me fix it, I'll be extrmrely thankful as the final code is due tonight. Thank you!...

Q:

It should be twi dismantions. and it should dyanamically create GTAs and TAs processes based on the the number of the chpaters (you can assume that each chapter will always have two assingments). The...

Q:

As provided, the code will draw and rotate a simple triangle in 3D. You are to modify the code to draw a cube in 3d with each vertex having its own color. Details: You are to draw a cube centered on...

Q:

As provided, the code will draw and rotate a simple triangle in 3D. You are to modify the code to draw a cube in 3d with each vertex having its own color. Details: You are to draw a cube centered on...

Q:

Please I want to modify the error with this code and run it, I want to run it after deleting the key function. Run in dev C++ and image the output #include #ifdef SUN #include #else #include #endif...

Q:

//Here is the correct code, you have defined the functions declaration at the end of code and function defination at the top which is why it was giving the compliation error //&&&&&&&&&&&&&&&&&&&&&&&...

Q:

For this part of the assignment you will write a pair of C++ template functions to read a series of items from an input file, and then print the items. Program Implement the following template...

Q:

C++ CSCI 241 Assignment 8, Part 1 Assignment For this part of the assignment you will write a pair of C++ template functions to read a series of items from an input file, and then print the items....

Q:

There are 3 parts, each one implements a new function to read a series of items from an input file. The first one: Program Implement the following template functions in a header file called sorts.h ....

Q:

The figure below shows the 60F water flow rates from the branches of a main supply line. Find the total pressure drop (PAPE) for soldered copper pipe. The LAB = 70 ft, LBC = 160 ft, and LcD = 100 ft....

Q:

1. Why is the WACC used in capital budgeting? 2. Explain some of the factors that affect the cost of capital and describe whether or not it is something that a company can control.

Q:

10 Work out the perimeter of this semicircle. Round your answer correct to 2 decimal places (2 d.p.). 24 cm

Q:

Seved Help 14 Wisconsin Snowmobile Corp. is considering a switch to level production Cost efficiencies would occur under level production, and aftertax costs would decline by $31,500, but inventory...

Q:

What tends to skew and distort Average Salaries in most Gender Pay Equity Studies?

Q:

The FedScope employment database has a number of Dimension Tables and a Single Fact Table, as shown in Table 7.1. Which columns/data elements in the Fact Table would be most useful in Pay Equity...

Q:

After Defining and Building a Multidimensional OLAP Cube, what is stored in the Cube?