== Team Members ==
# [mailto:mavillaflor@myseneca.ca?subject=GPU610 Mark Anthony Villaflor] (Leader; Wiki; Maze)# [mailto:hli206@myseneca.ca?subject=GPU610 Huachen Li] (Wiki; Prime Numbers)
# [mailto:ylei11@myseneca.ca?subject=GPU610 Yanhao Lei] (Wiki; Pi)
[mailto:mavillaflor@myseneca.ca;hli206@myseneca.ca;ylei11@myseneca.ca?subject=GPU610 eMail All]
0.00 3.84 0.00 1 0.00 0.00 _GLOBAL__sub_I__ZN9MazeDebugC2Ejj
0.00 3.84 0.00 1 0.00 0.00 _GLOBAL__sub_I_main
*Profiling with windows environment:
[[File:SDiagram.PNG]]
'''Summary:'''
=== Maze ===
[[File:Out2.png|thumb|Maze 3x3]]
[[File:Out1.png|thumb|Maze 10x10]]
'''Introduction:'''
The Maze program (by corzani@github) generates a maze image (of type PNG).
The following function creates the canvas of the Maze and initialize it by filling all the hexadecimal within each pixel. After the canvas is filled, the function calls the createImage function. Finally, the function write the result from the createImage function in a png file using the "lpng" extension.
void MazePng::toPng(unsigned int scale) {
...
color_type = PNG_COLOR_TYPE_RGB;
bit_depth = 8;
row_pointers = new png_bytep[height];
// This for loop filling all the hexadecimal within each pixel.
for (int i = 0; i < height; ++i) {
row_pointers[i] = new png_byte[width * 3];
// The inner for loop fills all three colours (red, green, blue) to the pixel.
for (int j = 0; j < width * 3; j += 3) {
row_pointers[i][j] = WALL;
row_pointers[i][j + 1] = WALL;
row_pointers[i][j + 2] = WALL;
}
}
createImage(row_pointers, 0);
...
}
The following function carves the path of the Maze by changing the colour of pixels to white.
void MazePng::createImage(png_bytep *row_pointers, unsigned int scale) {
...
//This loops sets the pixels of the row_pointers using the cells pattern
for (unsigned int y = 0; y < height; ++y) {
setPixel(row_pointers, 0, (y * 2) + 1, WALL);
for (unsigned int x = 0; x < width; ++x) {
switch ((cells[(y * width) + x] & 0xC0) >> 6) {
case 2:
temp[0] = PATH;
setPixel(row_pointers, 2 + (x * 2), (y * 2) + 1, temp[0]);
break;
case 1:
temp[1] = PATH;
setPixel(row_pointers, 1 + (x * 2), (y * 2) + 2, temp[1]);
break;
case 0:
setPixel(row_pointers, 2 + (x * 2), (y * 2) + 1, temp[0]);
setPixel(row_pointers, 1 + (x * 2), (y * 2) + 2, temp[1]);
break;
}
setPixel(row_pointers, 1 + (x * 2), (y * 2) + 1, PATH);
}
}
}
'''Testing Environment:'''
* Operating system: Windows 8.1 Enterprise 64-bit (6.3, Build 9600) (9600.winblue_ltsb.160812-0914)
12. In VS, go to menu Build -> Build Solution
== PHASE 3 =='''Analysis:'''
== Presentation ==[[File:SPDiagram.PNG]]
'''1. Introduction'''* MazeThe program generates a diagram shows that maze less than the size of 2,250,000 cells perform better in png filethe serial code. It takes 2 arguments: However, the height and the width of the mazeparallelized code has better performance if there are more cells to process.
* Maze image [[File:Out2.png|thumb|Maze 3x3]] [[File:Out1.png|thumb|Maze 10x10]] * Analysis: '''Parallelizable?''' The program function named toPng() has a runtime of O(n^2); it takes up an average of 42% (min: 35%; max: 50%) of the program's execution time.== PHASE 3 ==
for (int i = 0; i < height; ++i) { row_pointers[i] = new png_byte[width * 3]; for (int j = 0; j < width * 3; j +Attempt 1 = 3) { row_pointers[i][j] = WALL; row_pointers[i][j + 1] = WALL; row_pointers[i][j + 2] = WALL; } }
* Profiling:We tried to optimize the maze program by removing some of the if statements to reduce thread divergence, however the attempt made the program slower than the phase 2 version and only a little bit faster than the serial version.
[[File:SDiagram.PNG]] '''2. ParallelizeNew Kernels:''' *Original Code: /* * MazePng.cpp * * Created on: 12 Jul 2013 * Author: yac */ #include "MazePng.h" #include <stdio.h> #include <stdlib.h> #include <string.h> const png_byte WALL = 0x00; const png_byte PATH = 0xFF; MazePng::MazePng(const unsigned int width, const unsigned int height) : AbstractMaze(width, height) { } MazePng::~MazePng() { } void MazePng::toPng(unsigned int scale) { png_byte color_type; png_byte bit_depth; png_structp png_ptr; png_infop info_ptr; png_bytep *row_pointers; int height, width; FILE *fp; width = (this->width * 2) + 1; height = (this->height * 2) + 1; color_type = PNG_COLOR_TYPE_RGB; bit_depth = 8; row_pointers = new png_bytep[height]; for (int i = 0; i < height; ++i) { row_pointers[i] = new png_byte[width * 3]; for (int j = 0; j < width * 3; j += 3) { row_pointers[i][j] = WALL; row_pointers[i][j + 1] = WALL; row_pointers[i][j + 2] = WALL; } } createImage(row_pointers, 0); fp = fopen("out.png", "wb"); png_ptr = png_create_write_struct(PNG_LIBPNG_VER_STRING, NULL, NULL, NULL); info_ptr = png_create_info_struct(png_ptr); png_init_io(png_ptr, fp); png_set_IHDR(png_ptr, info_ptr, width, height, bit_depth, color_type, PNG_INTERLACE_NONE, PNG_COMPRESSION_TYPE_BASE, PNG_FILTER_TYPE_BASE); png_write_info(png_ptr, info_ptr); png_write_image(png_ptr, row_pointers); png_write_end(png_ptr, NULL); png_destroy_write_struct(&png_ptr, &info_ptr); for (int i = 0; i < height; ++i) { delete[] row_pointers[i]; } delete[] row_pointers; fclose(fp); } void inline MazePng::setPixel(png_bytep *row_pointers, unsigned int x, unsigned int y, png_byte type) { row_pointers[y][0 + 3 * x] = type; row_pointers[y][1 + 3 * x] = type; row_pointers[y][2 + 3 * x] = type; } void MazePng::createImage(png_bytep *row_pointers, unsigned int scale) { /* for (unsigned int x = 0; x < ((width * 2) + 1); ++x) { setPixel(row_pointers, x, 0, WALL); } */ if Initialize all pixels to black (start < widthhex 000) { setPixel(row_pointers, start * 2 + 1, 0, PATH); } png_byte temp[2] = { 0 }; for (unsigned int y = 0; y < height; ++y) { setPixel(row_pointers, 0, (y * 2) + 1, WALL); for (unsigned int x = 0; x < width; ++x) { switch ((cells[(y * width) + x] & 0xC0) >> 6) { case 2: temp[0] = PATH; setPixel(row_pointers, 2 + (x * 2), (y * 2) + 1, temp[0]); break; case 1: temp[1] = PATH; setPixel(row_pointers, 1 + (x * 2), (y * 2) + 2, temp[1]); break; case 0: setPixel(row_pointers, 2 + (x * 2), (y * 2) + 1, temp[0]); setPixel(row_pointers, 1 + (x * 2), (y * 2) + 2, temp[1]); break; } setPixel(row_pointers, 1 + (x * 2), (y * 2) + 1, PATH); } } } *Kernel: __global__ void k_drawWalls(png_byte* rows, const short* cells, const int width, const int height, const int len, const int size) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < size) { rows[i] = WALL; __syncthreads(); int px = i % len; int py = i / len; int x = (px - 1) / 2; int y = (py - 1) / 2; if (px > 0 && py > 0 && x < width && y < height) { int c = (cells[y * width + x] & 0xC0) >> 6; int idx = py * len + 3 * px; if (c == 2) { if (py % 2 > 0 && px % 2 == 0) { rows[idx] = rows[idx + 1] = rows[idx + 2] = PATH; } } else if (c == 1) { if (py % 2 == 0 && px % 2 > 0) { rows[idx] = rows[idx + 1] = rows[idx + 2] = PATH; } } else if (c == 0) { if ((py % 2 > 0 && px % 2 == 0) || (py % 2 == 0 && px % 2 > 0)) { rows[idx] = rows[idx + 1] = rows[idx + 2] = PATH; } } if (py % 2 > 0 && px % 2 > 0) { rows[idx] = rows[idx + 1] = rows[idx + 2] = PATH; } } } } *Profiling: [[File:SPDiagram.PNG]] *Sum up: **Maze size < 1500 * 1500**Maze size > 1500 * 1500 '''3. Optimize''' * New Kernel
__global__ void k_drawWalls(png_byte* rows, const short* cells, const int width, const int height, const int len, const int size) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
}
// Set pixels to white according to the pattern the cell belongs to
__global__ void k_drawPaths(png_byte* rows, const short* cells, const int width, const int height, const int len) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
}
Because there were too many else if condition in the old Kernel, we rewrite the Kernel to avoid divergent.'''Analysis:'''
[[File:SPODiagram.PNG]]
*Sum up: For this attempt, the kernel executes for each cell instead of for each byte (in phase 2); it is no longer processing more than 1 pixel in each cell for every thread, which may be the cause to the longer processing time. === Attempt 2 === **Still too many if statement We decided to use shared memory in the Kernel**Almost GPU to improve the speed of the program. However, the maze image does not show correctly: the same section paths were showing as randomly coloured pixels; the cause is due to the serialthreads setting only a part (1/3, 2/3) of a pixel's values to hexadecimal F.
