21
edits
Changes
→Assignment 3
Once we had the 1D array to store the image, we were able to start writing kernels for the compute-intensive functions of our processor. We created kernels for reflectImage(), enlargeImage() and negateImage() and began to test for any sort of performance increase.
For reference of what the image processor we are using actually does, we have attached the sample input and output image we have been using. This is the results for running the image processor without specifying an enlarge scale which will perform a negation and reflection vertically. [[File:input.jpg]] [[File:outputPGM.jpg]] Due to device limitations we were only able to profile our program up to an enlarged scale of 8 but our results still showed a performance increase as the enlarge scale got to 4. [[File:GpuA2Spreadsheet.png]] As seen from the results above , our parallel implementation of the image processor shows some significant performance increase as the enlarge scale gets larger. We also noticed that the c++ implementation of the processor seems to run at around the same time as the CUDA implementation ran faster when there was no enlarge and when the image is enlarge was only negated and reflected. However, once the image is scaled by any factor, there is a definite increase in performance from the CUDA implementation2. It i also worth noting We believe that this could be due to the profiled times form costly operation of cudaMemcpy() since the CUDA implementation seemed operations we are doing to vary a lot more than the c++ implementation which we think is from pixels are not that intensive, the variation in times time it takes for the cudaMemcpycudaMEmcpy() could easily start to runexceed the time of the transformations.
To continue our optimizations, we think that we could get more of a performance increase by minimizing the amount of data copied to and from the GPU. We are going to look into storing the image on the GPU until all transformations have been done and then copying the data back from the GPU.
=== Assignment 3 ===
For our optimizations to our parallel code we stuck to our plan and focused on reducing the amount of cudaMemcpy() used throughout each of the functions.
void Image::negateImage(Image& oldImage)
/*negates image*/
{
Image tempImage(N, M, Q);
int* d_temp = nullptr;
int* d_img = nullptr;
int size = N * M;
int nblocks = size / ntpb;
cudaMalloc((void**)&d_temp, size * sizeof(int));
cudaMalloc((void**)&d_img, size * sizeof(int));
cudaMemcpy(d_temp, tempImage.pixelVal, size * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_img, pixelVal, size * sizeof(int), cudaMemcpyHostToDevice);
negate << <nblocks, ntpb >> >(d_temp, d_img, size);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
printf("Error: %s\n", cudaGetErrorString(err));
cudaDeviceSynchronize();
cudaMemcpy(tempImage.pixelVal, d_temp, size * sizeof(int), cudaMemcpyDeviceToHost);
cudaFree(d_temp);
cudaFree(d_img);
oldImage = tempImage;
}
void Image::reflectImage(bool flag, Image& oldImage)
/*Reflects the Image based on users input*/
{
int rows = oldImage.N;
int cols = oldImage.M;
Image tempImage(oldImage);
int* d_temp = nullptr;
int* d_img = nullptr;
int size = rows * cols;
int nblocks = size / ntpb;
cudaMalloc((void**)&d_temp, size * sizeof(int));
cudaMalloc((void**)&d_img, size * sizeof(int));
cudaMemcpy(d_temp, tempImage.pixelVal, size * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_img, oldImage.pixelVal, size * sizeof(int), cudaMemcpyHostToDevice);
if (flag) {
horizontalReflect << <nblocks, ntpb >> >(d_temp, d_img, size, rows, cols);
}
else {
verticalReflect << <nblocks, ntpb >> >(d_temp, d_img, size, rows, cols);
}
cudaDeviceSynchronize();
cudaMemcpy(tempImage.pixelVal, d_temp, size * sizeof(int), cudaMemcpyDeviceToHost);
cudaFree(d_temp);
cudaFree(d_img);
oldImage = tempImage;
}
void Image::enlargeImage(int value, Image& oldImage)
/*enlarges Image and stores it in tempImage, resizes oldImage and stores the
larger image in oldImage*/
{
int rows, cols, gray;
int pixel;
int enlargeRow, enlargeCol;
rows = oldImage.N * value;
cols = oldImage.M * value;
gray = oldImage.Q;
Image tempImage(rows, cols, gray);
int r = oldImage.N;
int c = oldImage.M;
int* d_temp = nullptr;
int* d_img = nullptr;
int size = rows * cols;
int nblocks = size / ntpb;
cudaMalloc((void**)&d_temp, size * sizeof(int));
cudaMalloc((void**)&d_img, size * sizeof(int));
cudaMemcpy(d_temp, tempImage.pixelVal, size * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_img, oldImage.pixelVal, (r * c) * sizeof(int), cudaMemcpyHostToDevice);
enlarge << <nblocks, ntpb >> >(d_temp, d_img, size, value, c, cols);
cudaDeviceSynchronize();
//set the image's data
cudaMemcpy(tempImage.pixelVal, d_temp, size * sizeof(int), cudaMemcpyDeviceToHost);
//free device mem
cudaFree(d_temp);
cudaFree(d_img);
oldImage = tempImage;
}
As seen above, for all of the functions we converted to parallel, there is a lot of copying to and from the device. The use of a temp image makes this task even more intensive as we are basically copying the image to and from the device twice. In order to remove these cudaMemcpy() calls from each of the functions, we decided to simply create the image we are going to process directly onto the device allowing our parallel functions to operate on the pixels of the image without having to copy both the image and and a temporary image to and from the device. With this solution we were able to effectively reduce our 3 cudaMemcpy() calls per function from the parallel code into 0 cudaMemcpy() calls per function in our optimized code.
To implement this solution, we moved the intensive cudaMemcpy() calls to just be called when the image is created as seen below.
Image::Image(int numRows, int numCols, int grayLevels)
/* Creates an Image of numRows x numCols and creates the arrays for it*/
{
N = numRows;
M = numCols;
Q = grayLevels;
cudaMalloc((void**)&pixelVal, N * M * sizeof(int));
cudaMemset(pixelVal, 0, N * M * sizeof(int));
}
Image::Image(const Image& oldImage)
/*copies oldImage into new Image object*/
{
N = oldImage.N;
M = oldImage.M;
Q = oldImage.Q;
int sz = M * N * sizeof(int);
cudaMalloc((void**)&pixelVal, sz);
cudaMemcpy(pixelVal, oldImage.pixelVal, sz, cudaMemcpyDeviceToDevice);
}
void Image::operator=(const Image& oldImage)
/*copies oldImage into whatever you = it to*/
{
N = oldImage.N;
M = oldImage.M;
Q = oldImage.Q;
if (pixelVal) {
cudaFree(pixelVal);
}
int sz = M * N * sizeof(int);
cudaMalloc((void**)&pixelVal, sz);
cudaMemcpy(pixelVal, oldImage.pixelVal, sz, cudaMemcpyDeviceToDevice);
}
Since we decided to store the image directly onto the device, we also had to add some cudaMemcpy() calls to the functions that handled retrieving and delivering the image to the host. This process was originally handled in the getPixelVal() and setPixelVal() functions below.
int Image::getPixelVal(int row, int col)
/*returns the gray value of a specific pixel*/
{
return pixelVal[row * M + col];
}
void Image::setPixelVal(int row, int col, int value)
/*sets the gray value of a specific pixel*/
{
pixelVal[row * M + col] = value;
}
But now we need to transfer the data to and from the device which is where the image is stored. We created 2 new functions to handle this called getPixels() and setPixels().
void Image::getPixels(int row, int col, int sz, int* out) {
cudaMemcpy(out, pixelVal + (row * M + col), sz * sizeof(int), cudaMemcpyDeviceToHost);
}
void Image::setPixels(int row, int col, int sz, int* in) {
cudaMemcpy(pixelVal + (row * M + col), in, sz * sizeof(int), cudaMemcpyHostToDevice);
}
After making these changes we were able to greatly simplify our parallel functions as shown below.
void Image::negateImage(Image& oldImage)
/*negates image*/
{
Image tempImage(N, M, Q);
int* d_temp = tempImage.pixelVal;
int* d_img = pixelVal;
int size = N * M;
int nblocks = size / ntpb;
negate << <nblocks, ntpb >> >(d_temp, d_img, size);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
printf("Error: %s\n", cudaGetErrorString(err));
cudaDeviceSynchronize();
oldImage = tempImage;
}
void Image::reflectImage(bool flag, Image& oldImage)
/*Reflects the Image based on users input*/
{
int rows = oldImage.N;
int cols = oldImage.M;
Image tempImage(oldImage);
int* d_temp = tempImage.pixelVal;
int* d_img = oldImage.pixelVal;
int size = rows * cols;
int nblocks = size / ntpb;
if (flag) {
horizontalReflect << <nblocks, ntpb >> >(d_temp, d_img, size, rows, cols);
}
else {
verticalReflect << <nblocks, ntpb >> >(d_temp, d_img, size, rows, cols);
}
cudaDeviceSynchronize();
oldImage = tempImage;
}
void Image::enlargeImage(int value, Image& oldImage)
/*enlarges Image and stores it in tempImage, resizes oldImage and stores the
larger image in oldImage*/
{
int rows, cols, gray;
int pixel;
int enlargeRow, enlargeCol;
rows = oldImage.N * value;
cols = oldImage.M * value;
gray = oldImage.Q;
Image tempImage(rows, cols, gray);
int r = oldImage.N;
int c = oldImage.M;
int* d_temp = tempImage.pixelVal;
int* d_img = oldImage.pixelVal;
int size = rows * cols;
int nblocks = size / ntpb;
enlarge << <nblocks, ntpb >> >(d_temp, d_img, size, value, c, cols);
cudaDeviceSynchronize();
oldImage = tempImage;
}
Finally we profiled everything again to ensure that our optimizations would show a performance increase. We noticed a large decrease in the execution time after the optimizations and were satisfied with the outcome.
[[File:gpu optimized.png]]
The source code for our final assignment can be viewed on [https://github.com/mdbell/GPUAssignment GitHub]