From 7721132520b03907ef8437980393037c333f8799 Mon Sep 17 00:00:00 2001
From: hk <tjudjo@gmail.com>
Date: Wed, 11 Jan 2017 21:36:19 +0100
Subject: [PATCH] sprawko v2

---
 DergunPiotr-WaskoDominik/zad6/dok.tex | 83 +++++++++++++++++++++------
 1 file changed, 64 insertions(+), 19 deletions(-)

diff --git a/DergunPiotr-WaskoDominik/zad6/dok.tex b/DergunPiotr-WaskoDominik/zad6/dok.tex
index 600b895..cd48452 100644
--- a/DergunPiotr-WaskoDominik/zad6/dok.tex
+++ b/DergunPiotr-WaskoDominik/zad6/dok.tex
@@ -62,30 +62,59 @@ Grupa 1 & Piotr Dergun, Dominik Waśko & \ddmmyyyydate\today \tabularnewline
 
 \section*{Zadanie 6 - Rozmycie Gaussa GPU}
 
-Celem zadania jest wykonanie rozmycia obrazu wykorzystując algorytm Gaussa za pomocą karty graficznej oraz technologii CUDA. 
+Celem zadania jest wykonanie rozmycia obrazu wykorzystując algorytm Gaussa za pomocą karty graficznej oraz technologii CUDA.  Poniżej przedstawiono kod kernela, który oblicza rozmycie Gaussa. 
 \begin{lstlisting}
-__global__ void multiply(char *a,char *b,int *c, int N){
-	int i =blockIdx.x * blockDim.x + threadIdx.x;
-	int sum = 0;
-	int temp,temp2;
-	while(i<N*N){
-		temp = (int)floor((float)(i/N*N));
-		temp2 = i%N;
-		for(int j=0;j<N;j++){
-			sum +=a[temp + j] * b[j*N + temp2 ];
+__global__ void gauss(int rows, int cols, unsigned char *in_r, unsigned char *in_g, unsigned char *in_b,unsigned char *out_r, unsigned char *out_g, unsigned char *out_b){
+	const int r=5;
+	int ratio[r][r] ={  {1, 4, 7, 4, 1},
+					    {4, 16, 26, 16, 4},
+						{7, 26, 41, 26, 7},
+						{4, 16, 26, 16, 4},
+						{1, 4, 7, 4, 1}};
+
+	int N,i,row,col,poz,temp,sum_r,sum_g,sum_b,offset;
+	offset = 2;
+	N = rows * cols;
+	i = blockIdx.x * blockDim.x + threadIdx.x;
+	while(i<N){
+		col = i%cols;
+		row = (i/cols);
+		if(!(row<offset||row>rows-offset-1||col<offset||col>cols-offset-1)){
+			sum_r = 0;
+			sum_g = 0;
+			sum_b = 0;
+			row-=offset;
+			temp = col - offset;
+
+			for (int j=0; j<r; ++j,++row){
+				col= temp;
+				for (int k=0; k<r; ++k,++col){
+					poz = (row*cols) + col;
+					sum_r  += ratio[j][k] *  in_r[poz];
+					sum_g  += ratio[j][k] *  in_g[poz];
+					sum_b  += ratio[j][k] *  in_b[poz];
+				}
+			}
+			out_r[i] = sum_r / 273;
+			out_g[i] = sum_g / 273;
+			out_b[i] = sum_b / 273;
+		}else{
+			out_r[i] = in_r[i];
+			out_g[i] = in_g[i];
+			out_b[i] = in_b[i];
 		}
-		c[i]=sum;
-		sum =0;
 		i+=blockDim.x * gridDim.x;
 	}
-}	
+}
 \end{lstlisting}
- 
+
+Obraz do kernela przesyłany jest w postaci 3 tablic jednowymiarowych( po jednej dla każdego koloru). Konwersja z tablicy dwuwymiarowej A o rozmiarze NxM na jednowymiarową tablice B o rozmiarze N*M odbywa się wg wzoru: A[i][j] = B[M*i+j].
+  
 Zadanie zostało uruchomione na komputerze wyposażonym w karte graficzna NVIDIA GeForce GT 640M. Karta posiada dwa multiprocesory po 192 rdzenie CUDA, taktowane zegarem 709 MHz. Pamięć karty to 980MB. 
 
-Program testowano na tablicy o rozmiarze 2048x2048. Parametry dla jakich uruchomiono program to 3 rozmiary siatki: 16, 64, 256 oraz dla każdego rozmiaru siatki rozmiar bloku: 64,128,192,256,512,1024.
+Program testowano na obrazie o rozmiarze 24107x4491 pikseli. Parametry dla jakich uruchomiono program to 3 rozmiary siatki: 16, 64, 256 oraz dla każdego rozmiaru siatki rozmiar bloku: 64,128,192,256,320.
 
-Na rysunku 1. przedstawiono wykres zależności czasu obliczeń od rozmiaru bloku dla wszystkich trzech rozmiarów siatki. Można  zauważyć, że  czas obliczeń najbardziej spada przy zwiększeniu rozmiaru bloku do 192. Dalsze zwiększanie rozmiaru bloku nie powoduje ciągłego zwiększania szybkości, czas obliczeń utrzymuje się równym poziomie. Rozmiar siatki ma bardzo mały wpływ na szybkość obliczeń.  
+Na rysunku 1. przedstawiono wykres zależności czasu obliczeń od rozmiaru bloku dla wszystkich trzech rozmiarów siatki. Można  zauważyć, że  czas obliczeń najbardziej spada przy zwiększeniu rozmiaru bloku do 128. Dalsze zwiększanie rozmiaru bloku nie powoduje tak dużego zmniejszenia czasu obliczeń, a w pewnym momencie zaczyna powodować wzrost czasu . Również rozmiar siatki ma wpływ na szybkość obliczeń: im siatka większa tym bardziej maleje czas.   
 
 
 
@@ -97,7 +126,23 @@ Na rysunku 1. przedstawiono wykres zależności czasu obliczeń od rozmiaru blok
 \label{fig:figure2}
 \end{figure}
 
-
-
-
+\section*{Porównanie}
+Porównanie technologii OpenMP i MPI wykonano na komputerze MacPro3.1 wyposażonym w  2 procesory Xeon E5462, przy czym do dyspozycji było 6 rdzeni. Program z CUDA uruchomiono na innym komputerze wyposażonym w katę graficzną NVIDIA GeForce GT 640M z dwoma multiprocesorami po 192 rdzenie CUDA, taktowane zegarem 709 MHz. Do porównania użyto wykonanych na laboratoriach  programów obliczających rozmycie Gaussa . Do testów wybrano obraz o rozdzielczości 24107x4491. W tabeli przedstawiono czas wykonania przez poszczególne programy oraz liczbę użytych przez nie wątków, w przypadku CUDA jest to iloczyn rozmiaru siatki oraz rozmiaru bloku. 
+
+ \vspace{8mm} 
+\begin{table}
+\centering
+\begin{tabular}{|c|c|c|}
+  \hline 
+     & \textbf{wątki} & \textbf{czas[ms]}\\
+  \hline
+  \
+\textbf{OpenMP} & 6 & 2227\\
+  \hline
+  \textbf{MPI} & 6 & 1483\\
+  \hline
+  \textbf{CUDA} & 98304 & 571\\
+  \hline
+\end{tabular} 
+\end{table}
 \end{document}