==== Problem ====
After surveying the original code. We found one three major hot-spots for heavy CPU usage.
This block of code handles reshapes input pixels into a set of samples for classification.
const int N = width * height;[[File:SetSamplesSerial.png]] const int dim = imgThis block of code computes the distances between sampled centers and other input samples.channels(); cv[[File::Mat samples = cv::Mat(N, dim, CV_32FC1); for (int x = 0; x<width; x++) { for (int y = 0; y<height; y++) { for (int d = 0; d<dim; d++) {CalculateDistanceSerial.png|550px]] int index = y * width + x; samplesThis block of code generates the image that has to be outputted.at<float>(index, d) = (float)img.at<uchar>(y, x*dim + d); } } }[[File:GenerateImageSerial.png|550px]]
==== Analysis ====
You can find the new parallelized KmeansPlusPlus code
[https://github.com/agamdograMajinBui/KmeansPlusPlusCuda herekmeansplusplusCUDA].
Here is the kernel that we programmed.
__global__ void setCenter(float* d_center, float* d_sample, int n, int dim, int randi) { int i = blockIdxHere are the kernels that we programmed.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i < n && j < n) d_center[j * n + i] = d_sample[j * randi + i]; }
Launching the kernel
int nb = (n + ntpb - 1) / ntpb; dim3 dGrid(nb, nb, 1); dim3 dBlock(ntpb, ntpb, 1); float* d_center = nullptr; cudaMalloc((void**)&d_center, centers.rows * centers.cols * sizeof(float)); cudaMemcpy(d_center, (float*)centers.data, centers.rows * centers.cols * sizeof(float), cudaMemcpyHostToDevice); check(cudaGetLastError()); float* d_sample = nullptr; cudaMalloc((void**)&d_sample, samples.rows * samples.cols * sizeof(float)); cudaMemcpy(d_sample, (float*)samples.data, centers.rows * centers.cols * sizeof(float), cudaMemcpyHostToDevice); int rand = genrand_int31() % n; setCenter << <dGrid, dBlock >> >(d_center, d_sample, N, dim, rand); cudaDeviceSynchronize();Set Samples kernel
[[File:SetSamplesKernel.png|550px]]
The kernels:
setSamples - goes through the entire image and collects samples from the image (the current pixel and the next x number of pixels). __global__ void setSamples(cv::cuda::PtrStepSz<float> samples, cv::cuda::PtrStepSz<uchar> img, int dimC) { int i = blockIdx.y*blockDim.y + threadIdx.y; int j = blockIdx.x*blockDim.x + threadIdx.x; if (i >= img.rows || j >= img.cols) return; int index = i * img.cols + j; for (int d = 0; d<dimC; d++) { samples(index, d) = (float)img(i, j * dimC + d); } }Calculate Distance kernel
calculateDistance - goes through the image and computes the difference between the samples and the centers from the input image [[File:CalculateDistanceKernel.png|550px]]
__global__ void calculateDistance(cv::cuda::PtrStepSz<float> centers, cv::cuda::PtrStepSz<float> samples, int k, int N, int dim, double* minval, float* D) { // Compute distances between already sampled centers and other input samples. // Update nearest distance if it is smaller than previous ones. int col = blockIdx.x * blockDim.x + threadIdx.x; int row = blockIdx.y * blockDim.y + threadIdx.y; int i = col + row * N; //int i = blockIdx.y*blockDim.y + threadIdx.y; if (i >= N) return; double dist = 0.0; for (int d = 0; d<dim; d++) { double diff = centers(k - 1, d) - samples(i, d); dist += diff * diff; } if (dist < minval[i]) { minval[i] = dist; } *D += minval[i]; }Generate Image kernel
generateImage - takes the modified image and then writes it to the file using the function "out()" [[File:GenerateImageKernel.png|550px]]
__global__ void generateImage(cv::cuda::PtrStepSz<uchar> out, cv::cuda::PtrStepSz<int> indices, cv::cuda::PtrStepSz<float> centers, int dim) {
// Generate output image
int i = blockIdx.y*blockDim.y + threadIdx.y;
int j = blockIdx.x*blockDim.x + threadIdx.x;
if (i >= out.rows || j >= out.cols)
return;
int index = i * out.cols + j;
int ci = indices(index, 0);
for (int d = 0; d<dim; d++) {
out(i, j*dim + d) = (uchar)centers(ci, d);
}
}
After programming these kernel. we noticed an improvement in performace.
==== Conclusion ====
By comparing the run-times of the serial KmeansPlusPlus and the parallelized version, we can see that the performance of the program has improved slightly.
This program can further be improved by off-loading some more operations from the CPU to the GPU. But this will require more time and research[[File:GraphAssignment2.png|900px]]
The performance improvement is not significant for smaller clusters and iterations. But you can see that the performance has been improved for the higher test cases.
=== Assignment 3 ===
For assignment 3, we optimized the kernels by allocating the correct amounts of grids and block for each kernel. Previously, we allocated 32 threads by 32 blocks for every kernel call even when it did not require it. After adjustments, we found significant improvements for many of the kernels.
====Runtime of program====
Here, we see that the program was improved by the optimizations of threads per block.
Runtime of program:
For larger images, we found that the program was improved more and more as the amount of clusters and iterations increased.
[[File:Big Image.png]]
For medium images, we found more inconsistent results.
[[File:Med Image.png]]
For small images, we found the most inconsistent results after optimizations.
[[File:Small Image.png]]
When the image side increases, the more efficient the kernel.
====Runtime of each kernel====
Each kernel individually found significant or marginal improvements after adjusting for thread/block size.
Runtime of kernels:
Set samples found small improvements on average.
[[File:Set Samples.png]]
Here we changed the calculation of y_index to the outside of the inner loop.
[[File:SetSamplesKernelOptimized.png|550px]]
Calcuate distance found a significant improvements.
[[File:Calculate Distance Kernel.png]]
The biggest change was the thread/block size.
[[File:CalculateDistanceKernelOptimized.png|550px]]
Generate image found improvements as well since image sizes varied. Changing the thread/block size to the correct amount of pixels enabled better usage of memory.
[[File:Generate Image Kernel.png]]
The biggest change was the thread/block size.
[[File:GenerateImageKernelOptimized.png|550px]]
