From 106b2a76b22c5a0d1ed5a8d2282f92c51e7ca0bc Mon Sep 17 00:00:00 2001
From: Brandon Rodriguez <brodriguez8774@gmail.com>
Date: Sat, 24 Apr 2021 13:32:10 -0400
Subject: [PATCH] Correct CUDA thread generation and create color manipulation
test functions
---
part_1/image.cu | 139 +++++++++++++++++++++++++++++++++++++++++++-----
1 file changed, 126 insertions(+), 13 deletions(-)
diff --git a/part_1/image.cu b/part_1/image.cu
index 3b1c01d..87afa26 100644
--- a/part_1/image.cu
+++ b/part_1/image.cu
@@ -159,6 +159,120 @@ __global__ void cuda_alternating_blocks(unsigned long total_pixels, unsigned lon
}
}
+
+/**
+ * Updates image to increase percieved value of red in image.
+ * Purely for testing purposes.
+ */
+__global__ void cuda_increase_red(unsigned long total_pixels, unsigned long chunk_size, int chunk_counter, Image::PixelStruct* pixel_arr) {
+ // Calculate current index.
+ unsigned long pixel_index = blockIdx.x * blockDim.x + threadIdx.x;
+
+ // Adjust index for current chunk.
+ 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) {
+
+ // Calculate new pixel colors.
+ float red = pixel_arr[pixel_index].r + 0.1f;
+ float green = pixel_arr[pixel_index].g - 0.1f;
+ float blue = pixel_arr[pixel_index].b - 0.1f;
+
+ // Handle if colors exceed expected values.
+ if (red > 1.0f) {
+ red = 1.f;
+ }
+ if (green < 0.0f) {
+ green = 0.f;
+ }
+ if (blue < 0.0f) {
+ blue = 0.f;
+ }
+
+ // Update pixel values.
+ pixel_arr[pixel_index].r = red;
+ pixel_arr[pixel_index].g = green;
+ pixel_arr[pixel_index].b = blue;
+ }
+}
+
+
+/**
+ * Updates image to increase percieved value of green in image.
+ * Purely for testing purposes.
+ */
+__global__ void cuda_increase_green(unsigned long total_pixels, unsigned long chunk_size, int chunk_counter, Image::PixelStruct* pixel_arr) {
+ // Calculate current index.
+ unsigned long pixel_index = blockIdx.x * blockDim.x + threadIdx.x;
+
+ // Adjust index for current chunk.
+ 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) {
+
+ // Calculate new pixel colors.
+ float red = pixel_arr[pixel_index].r - 0.1f;
+ float green = pixel_arr[pixel_index].g + 0.1f;
+ float blue = pixel_arr[pixel_index].b - 0.1f;
+
+ // Handle if colors exceed expected values.
+ if (red < 0.0f) {
+ red = 0.f;
+ }
+ if (green > 1.0f) {
+ green = 1.f;
+ }
+ if (blue < 0.0f) {
+ blue = 0.f;
+ }
+
+ // Update pixel values.
+ pixel_arr[pixel_index].r = red;
+ pixel_arr[pixel_index].g = green;
+ pixel_arr[pixel_index].b = blue;
+ }
+}
+
+
+/**
+ * Updates image to increase percieved value of blue in image.
+ * Purely for testing purposes.
+ */
+__global__ void cuda_increase_blue(unsigned long total_pixels, unsigned long chunk_size, int chunk_counter, Image::PixelStruct* pixel_arr) {
+ // Calculate current index.
+ unsigned long pixel_index = blockIdx.x * blockDim.x + threadIdx.x;
+
+ // Adjust index for current chunk.
+ 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) {
+
+ // Calculate new pixel colors.
+ float red = pixel_arr[pixel_index].r - 0.1f;
+ float green = pixel_arr[pixel_index].g - 0.1f;
+ float blue = pixel_arr[pixel_index].b + 0.1f;
+
+ // Handle if colors exceed expected values.
+ if (red < 0.0f) {
+ red = 0.f;
+ }
+ if (green < 0.0f) {
+ green = 0.f;
+ }
+ if (blue > 1.0f) {
+ blue = 1.f;
+ }
+
+ // Update pixel values.
+ pixel_arr[pixel_index].r = red;
+ pixel_arr[pixel_index].g = green;
+ pixel_arr[pixel_index].b = blue;
+ }
+}
+
//endregion CUDA Methods.
@@ -457,13 +571,8 @@ void Image::compute() {
// We handle image in chunks, in case we're dealing with large images with more pixels than CUDA device blocks/threads.
int total_image_chunks = (total_pixels / cuda_chunk_size) + 1;
logger.info("Total Image Chunks: " + std::to_string(total_image_chunks));
-
- // Calculate desired blocks and threads.
- int cuda_device_desired_block_count = total_pixels / cuda_device_max_threads;
- if (cuda_device_desired_block_count <= 0) {
- cuda_device_desired_block_count = 1;
- }
- // TODO: Handle if desired block count exceeds device maximum count.
+ logger.info("Desired CUDA block count: " + std::to_string(cuda_device_max_blocks));
+ logger.info("Desired CUDA thread count: " + std::to_string(cuda_device_max_threads));
// Process image by iterating over all chunks.
for (int chunk_index = 0; chunk_index < total_image_chunks; chunk_index++) {
@@ -476,13 +585,17 @@ void Image::compute() {
cudaMemcpy(gpu_pixel_arr, pixel_arr, struct_byte_count, cudaMemcpyHostToDevice);
// Run GPU kernel logic.
- // cuda_alternating_pixels<<<cuda_device_desired_block_count, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, gpu_pixel_arr);
+ // cuda_alternating_pixels<<<cuda_device_max_blocks, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, gpu_pixel_arr);
- int pixel_block_size = width / 10;
- if (pixel_block_size <= 1) {
- pixel_block_size = 2;
- }
- cuda_alternating_blocks<<<cuda_device_desired_block_count, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, pixel_block_size, width, height, gpu_pixel_arr);
+ // int pixel_block_size = width / 10;
+ // if (pixel_block_size <= 1) {
+ // pixel_block_size = 2;
+ // }
+ // cuda_alternating_blocks<<<cuda_device_max_blocks, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, pixel_block_size, width, height, gpu_pixel_arr);
+
+ // cuda_increase_red<<<cuda_device_max_blocks, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, gpu_pixel_arr);
+ // cuda_increase_green<<<cuda_device_max_blocks, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, gpu_pixel_arr);
+ cuda_increase_blue<<<cuda_device_max_blocks, cuda_device_max_threads>>>(total_pixels, cuda_chunk_size, chunk_index, gpu_pixel_arr);
// Synchronize all GPU computations before iterating over next chunk.
cudaDeviceSynchronize();
--
GitLab