diff --git a/documents/references.md b/documents/references.md index a4ebe1197796da661c20e1001a4541cc8e1e527e..08e21a27b8135a6d1f28fd1d599bda33b5691d0a 100644 --- a/documents/references.md +++ b/documents/references.md @@ -86,3 +86,6 @@ Various references used in project. ### Portable Bitmap Formats <https://en.wikipedia.org/wiki/Netpbm> + +### Convolution Logic +* Gaussian Blur - <https://computergraphics.stackexchange.com/questions/39/how-is-gaussian-blur-implemented> diff --git a/part_1/image.cu b/part_1/image.cu index 38899be3c34b1d1c00ca061cff767b099ae3fec0..e81b6d2c2720cc764ebf2a7f85899051eb045a71 100644 --- a/part_1/image.cu +++ b/part_1/image.cu @@ -85,7 +85,7 @@ void display_cuda_device_properties() { __global__ void cuda_convolution_blur( unsigned long total_pixels, unsigned long chunk_size, int chunk_counter, int image_width, int image_height, int mask_width, int mask_height, - Image::PixelStruct* pixel_arr_orig, Image::PixelStruct* pixel_arr, Image::PixelStruct* mask_arr + Image::PixelStruct* pixel_arr_orig, Image::PixelStruct* pixel_arr, Image::PixelStruct* mask_arr, Image::PixelStruct mask_weight_struct ) { // Calculate current index. unsigned long pixel_index = blockIdx.x * blockDim.x + threadIdx.x; @@ -94,177 +94,81 @@ __global__ void cuda_convolution_blur( pixel_index += (chunk_size * chunk_counter); // Only proceed if GPU thread corresponds to a pixel in the image. Ignore all extra threads. - if (pixel_index <= total_pixels) { - // printf("total_pixels: %ld pixel_index: %ld\n", total_pixels, pixel_index); + if (pixel_index < total_pixels) { + + // if (pixel_index == 71) { + // printf("\n"); + // printf("total_pixels: %ld pixel_index: %ld image_width: %d image_height: %d\n", total_pixels, pixel_index, image_width, image_height); + // printf("\n"); + // for (int row_index = 0; row_index < image_height; row_index++) { + // for (int col_index = 0; col_index < image_width; col_index++) { + // unsigned long arr_index = col_index + (row_index * image_width); + + // printf(" [%d,%d] (%f) ", col_index, row_index, pixel_arr[arr_index].r); + // } + // printf("\n"); + // } + // printf("\n\n"); + // } // Calculate pixel location data. int image_row = (pixel_index / image_width) % image_height; int image_col = pixel_index % image_width; - int image_mid_row = image_height / 2; - int image_mid_col = image_width / 2; // Calculate mask location data. int mask_mid_row = mask_height / 2; int mask_mid_col = mask_width / 2; // Calculate pixel locations based on mask values. - int col_far_left = image_mid_col - mask_mid_col; - int col_far_right = image_mid_col + mask_mid_col; - int row_far_top = image_mid_row - mask_mid_row; - int row_far_bot = image_mid_row + mask_mid_row; - - // // For now, only compute on pixel in center of image. - // if ((image_row == image_mid_row) && (image_col == image_mid_col)) { - // printf("\n\n"); - // printf("total_pixels: %ld image_width: %d image_height: %d\n", total_pixels, image_width, image_height); - // printf("Middle pixel found: %ld image_col: %d image_row: %d\n", pixel_index, image_col, image_row); - // printf("image_middle_col: %d image_middle_row: %d mask_middle_col: %d mask_middle_row: %d\n", image_mid_col, image_mid_row, mask_mid_col, mask_mid_row); - - // printf("\n"); - // printf("Associated mask indexes:\n"); - // printf("Center: [%d,%d]\n", image_mid_col, image_mid_row); - // printf("CenterLeft: [%d,%d]\n", col_far_left, image_mid_row); - // printf("CenterRight: [%d,%d]\n", col_far_right, image_mid_row); - // printf("CenterTop: [%d,%d]\n", image_mid_col, row_far_top); - // printf("CenterBottom: [%d,%d]\n", image_mid_col, row_far_bot); - // printf("\n\n"); - - // // Display grid of mask values. - // for (int row_index = row_far_top; row_index <= row_far_bot; row_index++) { - // for (int col_index = col_far_left; col_index <= col_far_right; col_index++) { - // printf(" [%d,%d] ", col_index, row_index); - // } - // printf("\n"); - // } - // printf("\n\n"); - - // // Display grid of mask values. - // for (int row_index = 0; row_index < image_height; row_index++) { - // for (int col_index = 0; col_index < image_width; col_index++) { - // int a = col_index; - // int b = row_index; - // int c = a + (b * image_width); - // printf(" [%d,%d] (%d) ", a, b, c); - // } - // printf("\n"); - // } - // printf("\n\n"); - - // // Display grid of mask values. - // for (int row_index = row_far_top; row_index <= row_far_bot; row_index++) { - // for (int col_index = col_far_left; col_index <= col_far_right; col_index++) { - // int a = col_index; - // int b = row_index; - // int c = a + (b * image_width); - // printf(" [%d,%d] (%d) ", a, b, c); - // } - // printf("\n"); - // } - // printf("\n\n"); - - // pixel_arr[pixel_index].r = 1.f; - // pixel_arr[pixel_index].g = 0.f; - // pixel_arr[pixel_index].b = 0.f; - - // Overlay mask onto current pixel, and iterate through all associated pixels. - float convolution_total_r = 0.f; - float convolution_total_g = 0.f; - float convolution_total_b = 0.f; - float convolution_weight_total_r = 0.f; - float convolution_weight_total_g = 0.f; - float convolution_weight_total_b = 0.f; - int convolution_counter = 0; - int mask_index = 0; - for (int row_index = row_far_top; row_index <= row_far_bot; row_index++) { - for (int col_index = col_far_left; col_index <= col_far_right; col_index++) { - - // Double check that mask index is within image bounds. Ignore otherwise. - if ( ( (col_index >= 0) && (col_index < image_width) ) && ( (row_index >= 0) && (row_index < image_height) ) ) { - int arr_index = col_index + (row_index * image_width); - - // // // int a = col_index; - // // // int b = row_index; - // // // int c = a + (b * image_width); - // // // printf(" [%d,%d] (%d) (%f,%f,%f)", a, b, c, pixel_arr[mask_index].r, pixel_arr[mask_index].g, pixel_arr[mask_index].b); - - convolution_total_r += pixel_arr_orig[arr_index].r * mask_arr[mask_index].r; - convolution_total_g += pixel_arr_orig[arr_index].g * mask_arr[mask_index].g; - convolution_total_b += pixel_arr_orig[arr_index].b * mask_arr[mask_index].b; - convolution_counter += 1; - - - if ( (col_index == image_mid_col) && (row_index == image_mid_row) ) { - - // float mask_inverse = 1.f - mask_arr[mask_index].r; - // float updated_r = pixel_arr_orig[arr_index].r * mask_arr[mask_index].r; - // float updated_inv_r = pixel_arr_orig[pixel_index].r * mask_inverse; - // printf("[%d,%d] orig_r: %f mask_weight: %f mask_inv: %f\n", col_index, row_index, pixel_arr_orig[arr_index].r, mask_arr[mask_index].r, mask_inverse); - // printf("updated_r: %f updated_inv_r: %f\n", updated_r, updated_inv_r); - // updated_r = updated_r + updated_inv_r; - // // printf("total_pixels: %ld pixel_index: %ld mask_index: %d\n", total_pixels, pixel_index, mask_index); - // printf("updated_r: %f\n", updated_r); - // // printf("[%d,%d] orig_r: %f mask_weight: %f mask_inv: %f new_r: %f\n", col_index, row_index, pixel_arr_orig[arr_index].r, mask_arr[mask_index].r, mask_inverse, updated_r); - - - - } - - // Calculate updated red value. - float mask_inverse = 1.f - mask_arr[mask_index].r; - float updated_pix = pixel_arr_orig[arr_index].r * mask_arr[mask_index].r; - float updated_inv_pix = pixel_arr_orig[pixel_index].r * mask_inverse; - convolution_total_r = convolution_total_r + updated_pix + updated_inv_pix; - convolution_weight_total_r += mask_arr[mask_index].r; - - // Calculate updated green value. - mask_inverse = 1.f - mask_arr[mask_index].g; - updated_pix = pixel_arr_orig[arr_index].g * mask_arr[mask_index].g; - updated_inv_pix = pixel_arr_orig[pixel_index].g * mask_inverse; - convolution_total_g = convolution_total_g + updated_pix + updated_inv_pix; - convolution_weight_total_g += mask_arr[mask_index].g; - - // Calculate updated blue value. - mask_inverse = 1.f - mask_arr[mask_index].b; - updated_pix = pixel_arr_orig[arr_index].b * mask_arr[mask_index].b; - updated_inv_pix = pixel_arr_orig[pixel_index].b * mask_inverse; - convolution_total_b = convolution_total_b + updated_pix + updated_inv_pix; - convolution_weight_total_b += mask_arr[mask_index].b; - - - // // float red = pixel_arr_orig[arr_index].r * mask_arr[mask_index].r; - // // float green = pixel_arr_orig[arr_index].g * mask_arr[mask_index].g; - // // float blue = pixel_arr_orig[arr_index].b * mask_arr[mask_index].b; - // // pixel_arr[arr_index].r = red; - // // pixel_arr[arr_index].g = green; - // // pixel_arr[arr_index].b = blue; - } - mask_index += 1; + int col_far_left = image_col - mask_mid_col; + int col_far_right = image_col + mask_mid_col; + int row_far_top = image_row - mask_mid_row; + int row_far_bot = image_row + mask_mid_row; + + // Initialize for mask loop. + int mask_index = 0; + float convolution_total_r = 0.f; + float convolution_total_g = 0.f; + float convolution_total_b = 0.f; + + // Loop through all mask indexes and calculate weighted values. + for (int row_index = row_far_top; row_index <= row_far_bot; row_index++) { + for (int col_index = col_far_left; col_index <= col_far_right; col_index++) { + + // Double check that mask index is within image bounds. Ignore otherwise. + if ( ( (col_index >= 0) && (col_index < image_width) ) && ( (row_index >= 0) && (row_index < image_height) ) ) { + + unsigned long arr_index = col_index + (row_index * image_width); + convolution_total_r += (pixel_arr_orig[arr_index].r * mask_arr[mask_index].r); + convolution_total_g += (pixel_arr_orig[arr_index].g * mask_arr[mask_index].g); + convolution_total_b += (pixel_arr_orig[arr_index].b * mask_arr[mask_index].b); } - // printf("\n"); - } - // Average out values and apply to pixel index. - if (convolution_total_r > 0.f) { - pixel_arr[pixel_index].r = (convolution_total_r / convolution_counter); - // pixel_arr[pixel_index].r = (1.f / convolution_weight_total_r) * convolution_total_r; - } else { - pixel_arr[pixel_index].r = 0.f; - } - if (convolution_total_g > 0.f) { - pixel_arr[pixel_index].g = (convolution_total_g / convolution_counter); - // pixel_arr[pixel_index].g = (1.f / convolution_weight_total_g) * convolution_total_g; - } else { - pixel_arr[pixel_index].g = 0.f; - } - if (convolution_total_b > 0.f) { - pixel_arr[pixel_index].b = (convolution_total_b / convolution_counter); - // pixel_arr[pixel_index].b = (1 / convolution_weight_total_b) * convolution_total_b; - } else { - pixel_arr[pixel_index].b = 0.f; + mask_index += 1; } + } - // printf("\n\n"); - // } + // Calculate final pixel data based on mask loop values. + if (convolution_total_r > 0.f) { + // Handle for value greater than 0. + pixel_arr[pixel_index].r = (1.f / mask_weight_struct.r) * convolution_total_r; + } else { + pixel_arr[pixel_index].r = 0.f; + } + + if (convolution_total_g > 0.f) { + // Handle for value greater than 0. + pixel_arr[pixel_index].g = (1.f / mask_weight_struct.g) * convolution_total_g; + } else { + pixel_arr[pixel_index].g = 0.f; + } + + if (convolution_total_b > 0.f) { + // Handle for value greater than 0. + pixel_arr[pixel_index].b = (1 / mask_weight_struct.b) * convolution_total_b; + } else { + pixel_arr[pixel_index].b = 0.f; + } } } @@ -583,8 +487,8 @@ void Image::compute() { int total_image_chunks = (image_total_pixels / cuda_chunk_size) + 1; // Generate convolution mask dimensions. Should be either 10% of image size or 3 pixels large, whichever is greater. - int mask_width = width / 10; - int mask_height = height / 10; + int mask_width = width / 50; + int mask_height = height / 50; if ((mask_width % 2) == 0) { mask_width -= 1; } @@ -602,8 +506,10 @@ void Image::compute() { // Calculate additional mask initialization values. // Found via much struggling and trial + error. - float mask_col_adj = 1.f / (mask_middle_col + 1); - float mask_row_adj = 1.f / (mask_middle_row + 1); + // float mask_col_adj = 1.f / (mask_middle_col + 1); + // float mask_row_adj = 1.f / (mask_middle_row + 1); + float mask_col_adj = 1.f; + float mask_row_adj = 1.f; // printf("mask_middle_col: %d mask_middle_row: %d\n", mask_middle_col, mask_middle_row); // printf("mask_col_adj: %f mask_row_adj: %f\n", mask_col_adj, mask_row_adj); unsigned long mask_total_pixels = mask_width * mask_height; @@ -613,6 +519,7 @@ void Image::compute() { // printf("\n\n"); logger.info("Mask size: " + std::to_string(mask_width) + "x" + std::to_string(mask_height)); PixelStruct* mask_arr = new PixelStruct[mask_total_pixels]; + PixelStruct mask_weight_struct = PixelStruct(0.f); for (int mask_row_index = 0; mask_row_index < mask_height; mask_row_index++) { for (int mask_col_index = 0; mask_col_index < mask_width; mask_col_index++) { // Get current overall mask index, based on row and col values. @@ -632,12 +539,16 @@ void Image::compute() { mask_arr[mask_index].r = index_weight; mask_arr[mask_index].g = index_weight; mask_arr[mask_index].b = index_weight; + mask_weight_struct.r += index_weight; + mask_weight_struct.g += index_weight; + mask_weight_struct.b += index_weight; // mask_arr[mask_index].r = 1.f; // mask_arr[mask_index].g = 1.f; // mask_arr[mask_index].b = 1.f; } // printf("\n"); } + // printf("Total Weights: r: %f g: %f b: %f\n", mask_weight_struct.r, mask_weight_struct.b, mask_weight_struct.g); // printf("\n\n"); // Allocate memory on GPU. @@ -657,7 +568,7 @@ void Image::compute() { logger.info("Accounts for pixels " + std::to_string(cuda_chunk_size * chunk_index) + " through " + std::to_string((cuda_chunk_size * chunk_index) + cuda_chunk_size) ); // Run GPU logic to use convolution to create blur effect. - cuda_convolution_blur<<<cuda_device_max_blocks, cuda_device_max_threads>>>(image_total_pixels, cuda_chunk_size, chunk_index, width, height, mask_width, mask_height, gpu_pixel_arr_orig, gpu_pixel_arr, gpu_mask_arr); + cuda_convolution_blur<<<cuda_device_max_blocks, cuda_device_max_threads>>>(image_total_pixels, cuda_chunk_size, chunk_index, width, height, mask_width, mask_height, gpu_pixel_arr_orig, gpu_pixel_arr, gpu_mask_arr, mask_weight_struct); // Synchronize all GPU computations before iterating over next chunk. cudaDeviceSynchronize();