=== Boxblur on an image using opencv C++ Library (Max Fainshtein) Source Code for Vehicle Detection ===<syntaxhighlight lang="cpp">My suggested topic void detect_vehicles() { for the parallel program term project is an application that performs a boxblur on an image using open cv. This is made possible by using the opencv library installed through https://sourceforge.net/projects/opencvlibrary/(unsigned int i = 0; i < files/opencv-win/3.3.0size(); i++) { /opencv-3.3.0-vc14.exe/download or by using Visual Studios NuGet packages Load one image at the time and installing opencv.display it load_image(img, files[i]); win.native by Harry Y. Opencv was used to read images and access and modify the color of each pixel so that it is the average of the user defined box kernal. This application is running at Oset_image(n^2img) where n is the number of pixel rows and columns for the image. Running this program for images of various sizes resulted in the following:;

My topic is about Algorithms and the stress on the CPU and RAM of running them with large sizes of arrays. When using most algorithms with a set of arrays of a small size, the algorithms tend to finish faster then a second. When using these algorithms with larger sized arrays, based on my results it can be seen that the stress continues to increase as the size of the arrays increase, which in turn would also increase the execution time of the algorithms. In a situation when these algorithms are being called multiple times on large sized arrays, there would be an immense increase in execution time which may also result in the program becoming overall slow and/or hanging.

The project that has been chosen for For Assignment 2 is Max Fainshtein's Assignment 1 project, Boxblur on an image using opencv C++ Library and we have decided to implement parallelize the Box Blur algorithm because we realized that it would have been the only algorithm worth parallelizing using CUDA into .The main choice that made us decide to continue with Box Blur, as supposed of the Vehicle Detection program, is because it was the project only problem where we could createa kernel for and gain a lot of gain in order execution speed. The only way to eliminate optimize the stress on Vehicle Detection program was to enable CUDA for the CPUDLIB library, which it is not really in the scope of this assignment.<br><br>----

<br><br>Each block thread of the grid would represent a single pixel within the image that is being processed. <br><br>By implementing this kernel, the process time of the blur effect had made a significant improvement, compared to it's serial counter-part as illustrated in the results graph below. === Kernel Code ===<syntaxhighlight lang="cpp">__global__void blur(unsigned char* input_image, unsigned char* output_image, int width, int height, int neighbour) {  const unsigned int offset = blockIdx.x*blockDim.x + threadIdx.x; int x = offset % width; int y = (offset - x) / width; if (offset < width*height) {  float output_red = 0; float output_green = 0; float output_blue = 0; int hits = 0; for (int ox = -neighbour; ox < neighbour + 1; ++ox) { for (int oy = -neighbour; oy < neighbour + 1; ++oy) { if ((x + ox) > -1 && (x + ox) < width && (y + oy) > -1 && (y + oy) < height) { const int currentoffset = (offset + ox + oy*width) * 3; output_red += input_image[currentoffset]; output_green += input_image[currentoffset + 1]; output_blue += input_image[currentoffset + 2]; hits++; } } } output_image[offset * 3] = static_cast<unsigned char>(output_red / hits); output_image[offset * 3 + 1] = static_cast<unsigned char>(output_green / hits); output_image[offset * 3 + 2] = static_cast<unsigned char>(output_blue / hits); }}</syntaxhighlight>=== Launching the Kernel ===<syntaxhighlight lang="cpp">void filter(const Mat& input, Mat& output, int width, int height, int neighbour){ //Calculate total number of bytes of input and output image const int colorBytes = input.step * input.rows; const int grayBytes = output.step * output.rows;  unsigned char *d_input, *d_output;  //Allocate device memory cudaMalloc((void**)&d_input, width*height * 3 * sizeof(unsigned char)); cudaMalloc((void**)&d_output, width*height * 3 * sizeof(unsigned char));  //Copy data from OpenCV input image to device memory cudaMemcpy(d_input, input.ptr(), width*height * 3 * sizeof(unsigned char), cudaMemcpyHostToDevice); //cudaMemcpy(d_input, input.ptr(), colorBytes, cudaMemcpyHostToDevice);   dim3 blockDims(512, 1, 1); //Calculate grid size to cover the whole image  dim3 gridDims((unsigned int)ceil((double)(width*height * 3 / blockDims.x)), 1, 1); //Launch the color conversion kernel blur <br< <gridDims, blockDims >> >(d_input, d_output, input.cols, input.rows, neighbour);  //Synchronize to check for any kernel launch errors cudaDeviceSynchronize();  //Copy back data from destination device meory to OpenCV output image  cudaMemcpy(output.ptr(), d_output, width*height * 3 * sizeof(unsigned char), cudaMemcpyDeviceToHost); //Free the device memory cudaFree(d_input); cudaFree(d_output);}</syntaxhighlight>==== Graph ====[[File:boxFilterFirst.png | 750px]]

We had realized Upon using Nvidia's Visual Profiler it was evident that our implementation of a kernel had made we can make some massive improvements, compared to the serial version, but after profiling the Assignment 2 version we had noticed that we could still make improvementstry and improve our kernel even further.

The kernels had been executing concurrently but Nvidia's Visual Profiler showed that we were not using all the percentage of concurrency was quite lowStreaming Multi Processors to their maximum capability.

Initiate thread count based One way to address low compute utilization is to attempt to increase occupancy of each SM. According to Cuda's occupancy calculator the machine we were using for testing had a compute capability of 6.1. This means that each SM had 32 resident blocks and 2048 resident threads. To achieve maximum occupancy you would have 2048/32 = 64 threads/ block. To determine an appropriate grid size we would divide the total number of pixels by the 64 threads/block. This allows us to use dynamic grid sizing depending on Compute Capability the size of the CUDA deviceimage passed in.

The number of <syntaxhighlight lang="cpp>int iDevice;cudaDeviceProp prop;cudaGetDevice(&iDevice);cudaGetDeviceProperties(&prop, iDevice);int resident_threads = prop.maxThreadsPerMultiProcessor;int resident_blocks = 8;if (prop.major >= 3 && prop.major < 5) { resident_blocks = 16; }else if (prop.major >= 5 && prop.major <= 6) { resident_blocks = 32;}//determine threads that were initialized per /block had been calculated based on resident threads and blocksdim3 blockDims(resident_threads/resident_blocks,1,1); //Calculate grid size to cover the whole imagedim3 gridDims(pixels/blockDims.x);</syntaxhighlight> This resulted in a compute utilization increase from 33% to close 43% but unfortunately this did not yield much improvements.

==== Graph ====