equations they can be used to model and study magnetohydrodynamics. courtesy of wikipedia ("https://en.wikipedia.org/wiki/Navier%E2%80%93Stokes_equations")
=== problem Application Code to be parallelized===
The problem with this application comes in the main function trying to calculate the finite-difference
}
=== Initial Speed Tests ran with no optimization on linux ===
By using the command line argument cat /proc/cpuinfo
||12500 x 12500 || 220198||
|}
=== gprof ===
System Specifications
== Application 2 Calculating Pi==
This application is pretty straightforward, it calculates Pi to the decimal point which is given by the user. So an input of 10 vs 100,000 will calculate Pi to either the 10th or 100 thousandth decimal.
=== problem Application code to be parallelized ===
Inside the function calculate we have:
I Believe the 2 for loops will cause a delay in the program execution time.
=== Initial Speed Tests ran with no optimization on linux ===
for this test the linux VM has:
||500000 ||671163||
|}
=== gprof ===
'''
for (int i=0; i <= nx-1; i++)
{
if (i*dx >= 0.5 && i*dx <= 1)
u[i][it] = un[i][it-1] - c*dt/dx*(un[i][it-1]-un[i-1][it-1]);
}
}'''
u[k * nt + 0] = 1;
}
for (int it = 1; it <= nx - 1; it++)
{
u[m * nx + it] = un[m * nx + it - 1] - c*dt / dx*(un[m * nx + it - 1] - un[(m - 1) * nx + it - 1]);
}
}'''
The for loop is not needed.
=== INITIALIZE KERNEL ===
__global__ void Initalize(float* u, float* un, int nx, int nt, float dx)
{
}
=== CALCULATE WAVE KERNEL ===
This was the tricky part in converting the original code into the kernel.
The program takes 2 arrays. Let us say the X's represent the arrays below
Array __global__ void Calculate (float* u, float* un,int nx, int c, float dx, float dt) { int j = blockIdx.x * blockDim.x + threadIdx.x; int i = blockIdx.y * blockDim.y + threadIdx.y; // removes from instructions because no need to do this NX amount of times float total = c*dt / dx; if (i < nx && j < nx) { for (int it = 1; it <= nx- 1; it++) { if (i != 0 || i < nx ) { un[i * nx + it-1] = u[i * nx + it-1 Array 2]; __syncthreads(); u[it] = un[1 * nx + it - 1]; __syncthreads(); u[i * nx + it ] = un[i * nx + it- 1] - c*dt / dx* (un[i * nx + it - 1] - un[(i - 1) * nx + it - 1]); __syncthreads(); } } } ==== HOW THE ALGORITHM WORKS ==== This is focusing on the algorithm inside the CALCULATE Kernel only.
xxxxx xxxxx xxxxx xxxxx 1. We begin with 2 Arrays
xxxxx xxxxx [[File:2Arrazs.png]]
xxxxx xxxxx
xxxxx xxxxx 2. The first column of the First array is initialized by the INITIALIZE Kernel.
Upon initialization the 1st column of the first array gets set, this will be represented by o's[[File:Initialize.png]]
Array 1 Array 23. The second array copies the values from the first column of the First array
oxxxx xxxxx oxxxx xxxxx [[File:Copy1stColumn.png]]
oxxxx xxxxx 4. The First array copies a single value from the Second array
oxxxx xxxxx [[File:2ndCall.png]]
oxxxx xxxxx 5. The remaining values for the 2nd column of the First array are calculated through the Second array as follows.
The next kernel below will execute the following calucations[[File:3rdCall.png]]
1st: 6. Array 2 will copy The 2nd column of the first First array is now copied into the 2nd column of Array 1the Second array and the cycle is repeated until finished.
Array 1 Array 2[[File:LAstReset.png]]
oxxxx oxxxx oxxxx oxxxx == CPU VS GPU Loop Comparisons Only==
oxxxx oxxxx Executing the program again with a problem size of 2000 2000 or 4,000,000 we yield the following results.
oxxxx oxxxx Keep in mind these times are only for the kernel launches and not the program as a whole.
oxxxx oxxxx PARALLIZED GPU CODE Fist for loop - took - 0 millisecs 2nd for Loop - took - 0 millisecs Press any key to continue . . . ORIGINAL CPU CODE Initialize arrays loop - took - 17 milliseconds Fist for loop - took - 1 millisecs 2nd for Loop - took - 15373 millisecs Press any key to continue . . .
2nd: Array 1 will set the values in its first row to the values in Array 2's 1st column and 1st index.= OPTIMIZATION =
Array 1 Array 2== OVERALL EXECUTION OF PROGRAM FOR CPU, PARALLELIZED GPU AND OPTIMIZED CODE ==
o2xxx oxxxx oxxxx 2xxxx TIMES ARE IN MILLISECONDS
oxxxx oxxxx N Linux Visual No Parallel Parallized 2000 ^ 2 1160 20520 6749 5000 ^ 2 28787 127373 n/a 10000 ^ 2 124179 522576 n/a
oxxxx oxxxx === Windows Display Driver Crash for problem size > 2000 & 2000 ===
oxxxx oxxxx When I try to give the program an argument of over 2000 & 2000 it will inform me that the windows dispay driver has crashed and rebooted.
3rd: After some research I discovered that this is an issue caused by the kernel taking too long to execute. Finally Array 1 will calculate its 2nd column by performing Windows has a calculation on 2 columns in Array 2 which default time limit where it will be represented by reset the Thread Identifier as CUDA GPU if it thinks it is frozen due to the X dimension and the Y dimension being represented by the for loop iteratoramount of time it is taking to perform its calculations.
This is called the Timeout detection & recovery method (TDR). A potential solution I found on the CUDA programming forum on NVidea's website suggested I try the following in the registry:
Array 1 Array 2
o3xxx 33xxx To Change the Graphic device timeout, use the following steps. Exit all Apps and Programs. Press the WinKey+R keys to display the Run dialog. Type regedit.exe and click OK to open the registry editor. Navigate to the following registry key: HKEY_LOCAL_MACHINE\SYSTEM\CurrentControlSet\Control\GraphicsDrivers With the GraphicsDrivers key selected, on the Edit menu, click New, and then select the following registry value from the drop-down menu specific to your version of Windows (32 bit, or 64 bit):
o3xxx 33xxx (NOTE: The TdrDelay name is Case Sensitive) For 64 bit Windows a. Select QWORD (64-bit) value. b. Type TdrDelay as the Name and click Enter. c. Double-click TdrDelay and add 8 for the Value data and clickOK. The above potential solution did not solve my problem.... The second solution I found was to change one of the properties on the GPU device named: kernelExecTimeoutEnabled;This property supposedly controls whether or not the device can be timed out. A value of (1) means it can be timed out, while a value of (0) means it is disabled.
o3xxx 33xxx The above also did not solve my issue with the display driver crashing.
o3xxx 33xxx ==== Solution to Windows Display Driver Crashing ====
o3xxx 33xxx The best way to prevent this error from happening is to make sure the kernel does not take too long to execute... So I altered my code and switched the Kernel Launch statement from a 2D grid to a 1D grid.
This reduced the number of threads firing in the kernel. In the Calculate Kernel which is below you can see the old one had all the threads from the ( y dimension) sitting idle doing nothing except slowing down the execution.
==== PARALLELIZED CALCULATE WAVE KERNEL ====
__global__ void Calculate (float* u, float* un,int nx, int c, float dx, float dt)
{
int j = blockIdx.x * blockDim.x + threadIdx.x;
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i < nx && j < nx)
{
for (int it = 1; it <= nx- 1; it++)
{
if (i != 0 || i < nx )
{
un[i * nx + it-1] = u[i * nx + it-1];
__syncthreads();
u[it] = un[1 * nx + it - 1];
__syncthreads();
u[i * nx + it ] = un[i * nx + it- 1] - c * dt/dx * (un[i * nx + it - 1] - un[(i - 1) * nx + it - 1]);
__syncthreads();
}
}
}
===== OPTIMIZED CALCULATE WAVE KERNEL CHANGES=====
The code below has been altered to remove the (j) variable and combined the two (if) statements into one, so that we can reduce (Thread Divergence), as well as move the (- c*dt/dx* ) recurring instruction set, and place it into a variable called total, so that each thread is NOT performing the same operation which causes a decrease in performance.
// kernerl
__global__ void Calculate(float* u, float* un, int nx, int c, float dx, float dt)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
float total = c*dt / dx;
if (i < nx && i != 0)
{
for (int it = 1; it <= nx - 1; it++)
{
un[i * nx + it - 1] = u[i * nx + it - 1];
__syncthreads();
u[it] = un[1 * nx + it - 1];
__syncthreads();
u[i * nx + it] = un[i * nx + it - 1] - total * (un[i * nx + it - 1] - un[(i - 1) * nx + it - 1]);
__syncthreads();
}
}
}
With this optimized code it is now possible to execute with a problem size > 2000 & 2000.
==== ORIGINAL INITIALIZATION KERNEL ====
The Initialize kernel has also been redesigned. Below is the original:
__global__ void Initalize(float* u, float* un, int nx, int nt, float dx)
{
int j i = blockIdx.x * blockDim.x + threadIdx.x; int i j = blockIdx.y * blockDim.y + threadIdx.y; // removes from instructions because no need to do this NX amount of times float total = c*dt / dx; if (i < nx && j < nx) { for (int it = 1; it <= nx- 1; it++) { if (i !*dx >= 0 || .5 && i *dx < nx = 1) { un u[i * nx + it-1] = 2; __syncthreads(); } else { u[i * nx + it-] = 1]; __syncthreads(); } } } ===== OPTIMIZED INITIALIZATION KERNEL CHANGES ===== I removed the variable (j), removed the syncthreads() which were not needed, I also removed the function running on the CPU that initializes all indexes int he arrays to 0, and moved it into the GPU below. __global__ void Initalize(float* u[it] , float* un, int nx, int nt, float dx) { int i = un[1 blockIdx.x * nx blockDim.x + it - 1]threadIdx.x; __syncthreads if (i < nx) { for (int it = 0; it < nx;it++) { u[i * nx + it ] = 0; un[i * nx + it- 1] - c= 0; } if (i*dt / dx>= 0.5 && i* (undx <= 1) u[i * nx + it - 1] - un= 2; else u[(i * nx] = 1; } } == POST OPTIMIZATION - Execution Comparison Times== If you have not, please take a look at section 3.1.1.1(just above) * nx + , as it shows how the first iteration of optimization has been delivered. Below is a comparison of times from the original CPU to the newly optimized kernel execution. These comaprison times are for the WHOLE execution of the program, not just parts. These include memory transfers, allocation, de- 1]allocation and calculations. TIMES ARE IN MILLISECONDS N Linux Visual No Parallel Parallized Optimized_A (2000 ^ 2) 1160 | 20520 | 6749 | 971 (5000 ^ 2);28787 | 127373 | n/a | 1417 __syncthreads (10000 ^ 2); 124179 | 522576 | n/a | 3054 [[File:ParallelizedVSOptimized.png]] == SECOND OPTIMIZATION == === Shared Memory === In order to speed up the execution time I will incorporate shared data into the Calculate Kernel. The problem I am facing is determining in what way to use shared memory. } As I outlined above in section 2.2.2 regarding how to calculation on each Array is performed the program is calculating column by column and not rows by rows. However, it is also moving between rows after calculating each column. } }I can only allocate a static array and not dynamic so my shared memory will be the same size I use as my predefined ntpb variable, which represents the threads I use per block. So as of writing this, my ntpb variable is 32, therefor each shared array will be a size of 128 bytes.
I cannot copy the whole array into shared memory, and I cannot copy the array row by row, so we will also parallelize need to copy the Initialize function below:array column by column into shared memory.
As for the second array it has become clear that it is no longer needed, as we can simply use the shared memory array to perform the calculations of each column and save the results in the original arrays next column, then copy that column into the shared array and repeat the calculations.
=== SHARED MEMORY KERNEL ===
// kernerl __global__ void initCalculate(float* u, float* un, int nx, int c, float dx, float dt)
{
/ __shared__ float s[ntpb]; int i = blockIdx.x *Initialize all values to 0, switching from vector to arrayblockDim.x + threadIdx. Vector defaults to 0, array defaults to null*/x; for (int i t = threadIdx.x; float total = 0c*dt / dx; if (i < nx; && i++!= 0 && t != 0) { for (int k it = 01; k it < = nx- 1; kit++) { s[t - 1] = u[(i - 1) * nx +it - 1]; u[it] = s[1]; __syncthreads(); u[i * nx + kit] = 0s[t] - total * (s[t] - s[t - 1]); __syncthreads();
}
}
}
=== EXECUTION COMPARISON BETWEEN OPTIMIZED AND SHARED KERNELS ===
=== CPU VS GPU ===Below in milliseconds are the execution times for the former Kernel and new shared Kernel
Executing the program again with a problem size of {| class="wikitable sortable" border="1" cellpadding="5"|+ Time Comparison! n !! Optimized !! Shared |-||2000 x 2000 or 4,000,000 we yield the following results||971|| 661 |||-||5000 x 5000 ||1417|| 936 |||-||10000 x 10000 ||3054|| 2329 |||} == THIRD OPTIMIZATION ==
=== SAVING TRAVEL COSTS BY REMOVING THE UNNECESSARY ARRAY ===
Device As we discovered above, the second array is not necessary while we are performing all the calculations on Shared Memory which can be seen in section 3.3.2. This provides us with compute capability 3the ability to further optimize our Kernel by reducing the amount of time we spend transferring data across the PCI bus. Below is an image of the data transfer times for the CALCULATE kernel.0 found (index 0)
Name: Quadro K2000
Compute Capability: 3Since both of the original Arrays are not needed in the final Kernel solution, we can save 50% of our transfer time across the PCI bus by removing one of the arrays.0
Total Global Memory: 2147483648
Max Threads per block[[File: 1024MEmCpy10000.png]]
maxGridSize: 002EF650=== GETTING 100% OCCUPANCY PER MULTIPROCESSOR===
maxThreadsDim: 002EF644'''Occupancy Calculator
Clock Rate in khz: 954000The CUDA Toolkit includes a spreadsheet that accepts as parameters the compute capability, the number of threads per block, the number of registers per thread and the shared memory per block. This spreadsheet evaluates these parameters against the resource limitations of the specified compute capability. This spreadsheet is named CUDA_Occupancy_Calculator.xls and stored under the ../tools/ sub-directory of the installed Toolkit.'''
Source--> https://scs.senecac.on.ca/~gpu610/pages/content/resou.html
With the existing CALCULATE Kernel the CUDA Occupancy Calculator is providing the following statistics as shown below... [[File:OriginalCalculator.png]] The current CALCULATE Kernel is only utilizing 50% of the MultiProcessor as shown above. Fist If the threads per block are switched from 32 to 512 we will achieve 100% occupancy as shown below. [[File:100Calculator.png]] === CALCULATE KERNEL === Here is the final CALCULATE Kernel for loop the application.The changes include removal of the second array. // kernerl __global__ void Calculate(float* u, int nx, int c, float dx, float dt) { __shared__ float s[ntpb]; int i = blockIdx.x * blockDim.x + threadIdx.x; int t = threadIdx.x; float total = c*dt / dx; if (i < nx && i != 0 && t != 0) { for (int it = 1; it <= nx - 1; it++) { s[t - took 1] = u[(i - 0 millisecs1) * nx + it - 1]; u[it] = s[1]; __syncthreads(); 2nd for Loop u[i * nx + it] = s[t] - total * (s[t] - took s[t - 0 millisecs1]); __syncthreads(); } } } === OPTIMIZATION TIME COMPARISONS === Below is a graph comparing times between Optimizations illustrating the amount of execution time saved in each iteration. The times are listed in milliseconds. [[File:OPTIMIZATIONCOMPARISON.png]] = CONCLUSIONS = == OVERALL TIME COMPARISONS == Below are the final comparisons of all execution times between the CPU and GPU. All times are in milliseconds. [[File:finalCompare.png]] == APPLICATION OUTPUT == Upon completion of the application it will create a file based on the output of the algorithm. Press any key The following image below displays that output comparing the original program to continue the parallelized program. [[File:outputs. png]] == FINAL THOUGHTS == Upon completion of this Project I have learned a few things: First, I learned that not all program can be parallelized even if they seem to be a good candidate to begin with.
As compared Secondly, understand the algorithm of the application is a key factor in being able to optimize the solution, because sometimes you will need to rearrange the CPU version code in order to obtain better performance from the GPU and understanding the algorithm will help ensure that the output at the end of the original programwill remain the same.
Fist for loop Thirdly the management of resources and constraints, having registers, shared memory, constant memory, latency, threads, and multi- took - 1 millisecsprocessors are all factors which need to be considered when using the GPU. 2nd for Loop - took - 15373 millisecs Press any key Understanding how these resources can impact and influence your program helps deciding which ones to continue . . use in specific situations.