MAJ

[kahina_paper1.git] / Root.tex

diff --git a/Root.tex b/Root.tex
index c845ba8162f8bc0e5ab621f14ef5629ef788d522..b4edff1c6bb4050f04aabdd4647293ddbb513e4c 100644 (file)
--- a/Root.tex
+++ b/Root.tex
@@ -425,9 +425,9 @@ Let Nbr\_iter the number of iteration necessary to compute all the roots,so the
 \begin{equation}\r
 Total\_time_{exe}=\left[N\left(T_{i}(N)+T_{j}\right)+O(n)\right].Nbr\_iter.\r
 \end{equation}\r
 \begin{equation}\r
 Total\_time_{exe}=\left[N\left(T_{i}(N)+T_{j}\right)+O(n)\right].Nbr\_iter.\r
 \end{equation}\r

-The execution time increase with the increasing of the polynomial's root, which take necessary to parallilize this step to reduce the execution time. In the following paper you explain how we parrallelize  this step using GPU architecture with CUDA platform.\r
+The execution time increase with the increasing of the polynomial's root, which take necessary to parallelize this step to reduce the execution time. In the following paper you explain how we parrallelize  this step using GPU architecture with CUDA platform.\r

 \r
 \r

-\subsubsection{Parralelize the steps on GPU }\r
+\subsubsection{Parallelize the steps on GPU }\r

 On the CPU Aberth algorithm both steps 3 and 4 contain the loop \verb=for= , it use one thread to execute all the instruction in the loop N times.Here we explain how the GPU architecture can compute this loop and reduce the execution time.\r
 The GPU architecture affect the execution of this loop to a groups of parallel threads organized as a grid of blocks each block contain a number of threads. All threads within a block are executed concurrently in parallel. the instruction are executed as a kernel.\r
 \r
 On the CPU Aberth algorithm both steps 3 and 4 contain the loop \verb=for= , it use one thread to execute all the instruction in the loop N times.Here we explain how the GPU architecture can compute this loop and reduce the execution time.\r
 The GPU architecture affect the execution of this loop to a groups of parallel threads organized as a grid of blocks each block contain a number of threads. All threads within a block are executed concurrently in parallel. the instruction are executed as a kernel.\r
 \r


@@ -440,7 +440,7 @@ Total\_time_{exe}=\left[\frac{N}{nbr\_thread}\left(T_{i}(N)+T_{j}\right)+O(n)\ri
 In theory, the $Total\_time_{exe}$ on GPU is speed up nbr\_thread times as a $Total\_time_{exe}$ on CPU. We show more details in the experiment part. \r
 ~\\\r
 ~\\\r
 In theory, the $Total\_time_{exe}$ on GPU is speed up nbr\_thread times as a $Total\_time_{exe}$ on CPU. We show more details in the experiment part. \r
 ~\\\r
 ~\\\r

-In CUDA platform, All the instruction of the loop \verb=for= are executed by the GPU as a kernel form. A kernel is a procedure written in CUDA and defined by a heading \verb=__global__=, which means that it is to be executed by the GPU.the following algorithm see the kernels created for the step 3 and 4 of the Aberth algorithm:\r
+In CUDA platform, All the instruction of the loop \verb=for= are executed by the GPU as a kernel form. A kernel is a procedure written in CUDA and defined by a heading \verb=__global__=, which means that it is to be executed by the GPU.the following algorithm see the Aberth algorithm on GPU:\r

 \r
 \begin{algorithm}[H]\r
 \LinesNumbered\r
 \r
 \begin{algorithm}[H]\r
 \LinesNumbered\r


@@ -455,20 +455,43 @@ tolerance threshold),P(Polynomial to solve)}
 \r
 Initialization of the parameter of the polynomial to solve\;\r
 Initialization of the solution vector $Z^{0}$\;\r
 \r
 Initialization of the parameter of the polynomial to solve\;\r
 Initialization of the solution vector $Z^{0}$\;\r

+Allocate and fill the data in the global memory GPU\;\r

 \r
 \While {$\Delta z_{max}\succ \epsilon$}{\r
  Let $\Delta z_{max}=0$\;\r
 \r
 \While {$\Delta z_{max}\succ \epsilon$}{\r
  Let $\Delta z_{max}=0$\;\r

-$\prec DimGrid,DimBloc\succ kernel\_save(d\_Z^{k-1});$\r
-$\prec DimGrid,DimBloc \succ kernel\_update(d\_z^{k});$\r
-\r
-$\prec DimGrid,DimBloc\succ kernel\_testConverge (d_?z_{max},d_Z^{k},d_Z^{k-1});$\r
+$ kernel\_save(d\_Z^{k-1})$\;\r
+$ kernel\_update(d\_z^{k})$\;\r
+$kernel\_testConverge (d_?z_{max},d_Z^{k},d_Z^{k-1})$\;\r

 }\r
 \end{algorithm}\r
 ~\\ \r
 \r
 }\r
 \end{algorithm}\r
 ~\\ \r
 \r

+After the initialization step, all data of the root finding problem to be solved must be copied from the CPU memory to the GPU global memory, because the GPUs only work on the data filled in their memories. Next, all the data-parallel arithmetic operations inside the main loop \verb=(do ... while(...))= are executed as kernels by the GPU. The first kernel \textit{save} in line( 6, Algorithm 2) consist to save the vector of polynomial's root found at the previous time step on GPU memory, in order to test the convergence of the root at each iteration in line (8, Algorithme2).\r
+\r
+The second kernel executes the iterative function and update Z(k),as formula (), we notice that the kernel update are called in two forms,  separated with  the value of R which determines the radius beyond which we apply the logarithm formula like this: \r
+\r
+\begin{algorithm}[H]\r
+\LinesNumbered\r
+\caption{A global Algorithm for the iterative function}\r
+\r
+\eIf{$(\left|Z^{(k)}\right|<= R)$}{\r
+$kernel\_update(d\_z^{k})$\;}\r
+{\r
+$kernel\_update\_Log(d\_z^{k})$\;\r
+}\r
+\end{algorithm}\r
+\r
+The first form execute the formula(8) if all the module's $( |Z(k)|<= R)$, else the kernel execute the formulas(13,14).the radius R was computed like:\r
+\r
+$$R = \exp( \log(DBL\_MAX) / (2*(double).N) )$$\r
+\r
+where N the degree of the polynomial,DBL\_MAX is the maximum value of a double.  \r
+The last kernel verify the convergence of the root after each update of $Z^{(k)}$, as formula(), we used the function of the CUBLAS Library (CUDA Basic Linear Algebra Subroutines) to implement this kernel. \r
+\r
+The kernels terminates its computations when all the root are converged. Finally, the solution of the  root finding problem is copied back from the GPU global memory to the CPU memory. We use the communication functions of CUDA for the memory allocations in the GPU \verb=(cudaMalloc())= and the data transfers from the CPU memory to the GPU memory \verb=(cudaMemcpyHostToDevice)=\r
+or from the GPU memory to the CPU memory \verb=(cudaMemcpyDeviceToHost))=. \r

 \r
 \r
 \r
 \r

-here we need to create two kernel for the step 3 \textit{Kernel\_save} is used to save vector $Z^{K-1}$ and \textit{kernel\_update} is used to update the $Z^{k}$ vector. In phase 4 a kernel is created to test the convergence of the method\r

 \r
 \r
 \subsubsection{the kernel corresponding }\r
 \r
 \r
 \subsubsection{the kernel corresponding }\r