revu s4.B

[kahina_paper2.git] / paper.tex

diff --git a/paper.tex b/paper.tex
index 2f436392e9e952cd06c8598e1b38db6ae745dc9c..4f0138ec95a1b34f3b1b87a8ad1e6b3980739e60 100644 (file)
--- a/paper.tex
+++ b/paper.tex
@@ -463,7 +463,7 @@ Open Multi-Processing (OpenMP) is a shared memory architecture API that provides
 a portable approach for parallel programming on shared memory systems based on compiler directives, that can be included in order
 to parallelize a loop. In this way, a set of loops can be distributed along the different threads that will access to different data allocated in local shared memory. One of the advantages of OpenMP is its global view of application memory address space that allows relatively fast development of parallel applications with easier maintenance. However, it is often difficult to get high rates of performance in large scale applications. Although usage of OpenMP  threads and managed data explicitly done with MPI can be considered, this approcache undermines the advantages of OpenMP.
 
 a portable approach for parallel programming on shared memory systems based on compiler directives, that can be included in order
 to parallelize a loop. In this way, a set of loops can be distributed along the different threads that will access to different data allocated in local shared memory. One of the advantages of OpenMP is its global view of application memory address space that allows relatively fast development of parallel applications with easier maintenance. However, it is often difficult to get high rates of performance in large scale applications. Although usage of OpenMP  threads and managed data explicitly done with MPI can be considered, this approcache undermines the advantages of OpenMP.
 

-%\subsection{OpenMP} %L'article en Français Programmation multiGPU – OpenMP versus MPI
+%\subsection{OpenMP} 

 %OpenMP is a shared memory programming API based on threads from
 %the same system process. Designed for multiprocessor shared memory UMA or
 %NUMA [10], it relies on the execution model SPMD ( Single Program, Multiple Data Stream )
 %OpenMP is a shared memory programming API based on threads from
 %the same system process. Designed for multiprocessor shared memory UMA or
 %NUMA [10], it relies on the execution model SPMD ( Single Program, Multiple Data Stream )


@@ -474,10 +474,10 @@ to parallelize a loop. In this way, a set of loops can be distributed along the
 %have private memory areas [6].
 
 \subsection{MPI} 
 %have private memory areas [6].
 
 \subsection{MPI} 

-The MPI (Message Passing Interface) library allows to create computer programs that run on a distributed memory architecture. The various processes have their own environment of execution and execute their code in a asynchronous way, according to the MIMD model  (Multiple Instruction streams, Multiple Data streams); they communicate and synchronise by exchanging messages~\cite{Peter96}. MPI messages are explicitly sent, while the exchanges are implicit within the framework of a multi-thread programming environment like OpenMP or Pthreads.
+The MPI (Message Passing Interface) library allows to create computer programs that run on a distributed memory architecture. The various processes have their own environment of execution and execute their code in a asynchronous way, according to the MIMD model  (Multiple Instruction streams, Multiple Data streams); they communicate and synchronize by exchanging messages~\cite{Peter96}. MPI messages are explicitly sent, while the exchanges are implicit within the framework of a multi-thread programming environment like OpenMP or Pthreads.

  
  

-\subsection{CUDA}%L'article en anglais Multi-GPU and multi-CPU accelerated FDTD scheme for vibroacoustic applications
-CUDA (an acronym for Compute Unified Device Architecture) is a parallel computing architecture developed by NVIDIA~\cite{NVIDIA12}. The
+\subsection{CUDA}
+CUDA (an acronym for Compute Unified Device Architecture) is a parallel computing architecture developed by NVIDIA~\cite{CUDA10}. The

 unit of execution in CUDA is called a thread. Each thread executes a kernel by the streaming processors in parallel. In CUDA,
 a group of threads that are executed together is called a thread block, and the computational grid consists of a grid of thread
 blocks. Additionally, a thread block can use the shared memory on a single multiprocessor while the grid executes a single
 unit of execution in CUDA is called a thread. Each thread executes a kernel by the streaming processors in parallel. In CUDA,
 a group of threads that are executed together is called a thread block, and the computational grid consists of a grid of thread
 blocks. Additionally, a thread block can use the shared memory on a single multiprocessor while the grid executes a single


@@ -487,9 +487,9 @@ scope. The effective bandwidth of each memory space depends on the memory access
 bandwidth than the shared memory, the global memory accesses should be minimized.
 
 
 bandwidth than the shared memory, the global memory accesses should be minimized.
 
 

-We introduced three paradigms of parallel programming. Our objective consist to implement an algorithm of root finding polynomial on multiple GPUs. It primordial to know how to manage CUDA contexts of different GPUs. A direct method for controlling the various GPU is to use as many threads or processes as GPU devices. We can choose the GPU index based on the identifier of OpenMP thread or the rank of the MPI process. Both approaches will be investigated.
+We introduced three paradigms of parallel programming. Our objective consists in implementing a root finding polynomial algorithm on multiple GPUs. To this end, it is primordial to know how to manage CUDA contexts of different GPUs. A direct method for controlling the various GPUs is to use as many threads or processes as GPU devices. We can choose the GPU index based on the identifier of OpenMP thread or the rank of the MPI process. Both approaches will be investigated.

 
 

-\section{The EA algorithm on single GPU}
+\section{The EA algorithm on a single GPU}

 \subsection{the EA method}
 
 A cubically convergent iteration method to find zeros of
 \subsection{the EA method}
 
 A cubically convergent iteration method to find zeros of


@@ -536,7 +536,7 @@ As for any iterative method, we need to choose $n$ initial guess points $z^{0}_{
 The initial guess is very important since the number of steps needed by the iterative method to reach
 a given approximation strongly depends on it.
 In~\cite{Aberth73} the Ehrlich-Aberth iteration is started by selecting $n$
 The initial guess is very important since the number of steps needed by the iterative method to reach
 a given approximation strongly depends on it.
 In~\cite{Aberth73} the Ehrlich-Aberth iteration is started by selecting $n$

-equi-spaced points on a circle of center 0 and radius r, where r is
+equi-distant points on a circle of center 0 and radius r, where r is

 an upper bound to the moduli of the zeros. Later, Bini and al.~\cite{Bini96}
 performed this choice by selecting complex numbers along different
 circles which relies on the result of~\cite{Ostrowski41}.
 an upper bound to the moduli of the zeros. Later, Bini and al.~\cite{Bini96}
 performed this choice by selecting complex numbers along different
 circles which relies on the result of~\cite{Ostrowski41}.


@@ -555,9 +555,9 @@ v_{i}=\frac{|\frac{a_{n}}{a_{i}}|^{\frac{1}{n-i}}}{2}.
 \end{equation}
 
 \subsubsection{Iterative Function}
 \end{equation}
 
 \subsubsection{Iterative Function}

-The operator used by the Aberth method is corresponding to the
-following equation~\ref{Eq:EA1} which will enable the convergence towards
-polynomial solutions, provided all the roots are distinct.
+The operator used by the Aberth method  corresponds to the
+equation~\ref{Eq:EA1}, it enables the convergence towards
+the polynomials zeros, provided all the roots are distinct.

 
 %Here we give a second form of the iterative function used by the Ehrlich-Aberth method: 
 
 
 %Here we give a second form of the iterative function used by the Ehrlich-Aberth method: 
 


@@ -597,14 +597,14 @@ to execute in parallel must be made by the GPU to accelerate
 the execution of the application, like the step 3 and step 4. On the other hand, all the
 sequential operations and the operations that have data
 dependencies between threads or recursive computations must
 the execution of the application, like the step 3 and step 4. On the other hand, all the
 sequential operations and the operations that have data
 dependencies between threads or recursive computations must

-be executed by only one CUDA or CPU thread (step 1 and step 2). Initially we specifies the organization of threads in parallel, need to specify the dimension of the grid Dimgrid: the number of block per grid and block by DimBlock: the number of threads per block required to process a certain task. 
-
-we create the kernel, for step 3 we have two kernels, the
-first named \textit{save} is used to save vector $Z^{K-1}$ and the kernel
-\textit{update} is used to update the $Z^{K}$ vector. In step 4 a kernel is
-created to test the convergence of the method. In order to
-compute function H, we have two possibilities: either to use
-the Jacobi method, or the Gauss-Seidel method which uses the
+be executed by only one CUDA or CPU thread (step 1 and step 2). Initially, we specify the organization of parallel threads, by specifying the dimension of the grid Dimgrid, the number of blocks per grid DimBlock and the number of threads per block. 
+
+The code is organzed  by what is named kernels, portions o code that are run on GPU devices. For step 3, there are two kernels, the
+first named \textit{save} is used to save vector $Z^{K-1}$ and the seconde one is named
+\textit{update} and is used to update the $Z^{K}$ vector. For step 4, a kernel 
+tests the convergence of the method. In order to
+compute the function H, we have two possibilities: either to use
+the Jacobi mode, or the Gauss-Seidel mode of iterating which uses the

 most recent computed roots. It is well known that the Gauss-
 Seidel mode converges more quickly. So, we used the Gauss-Seidel mode of iteration. To
 parallelize the code, we created kernels and many functions to
 most recent computed roots. It is well known that the Gauss-
 Seidel mode converges more quickly. So, we used the Gauss-Seidel mode of iteration. To
 parallelize the code, we created kernels and many functions to


@@ -618,9 +618,7 @@ implement, as the development of corresponding kernels with
 CUDA is longer than on a CPU host. This comes in particular
 from the fact that it is very difficult to debug CUDA running
 threads like threads on a CPU host. In the following paragraph
 CUDA is longer than on a CPU host. This comes in particular
 from the fact that it is very difficult to debug CUDA running
 threads like threads on a CPU host. In the following paragraph

-Algorithm 1 shows the GPU parallel implementation of Ehrlich-Aberth method.
-
-Algorithm~\ref{alg2-cuda} shows a sketch of the Ehrlich-Aberth method using CUDA.
+Algorithm~\ref{alg1-cuda} shows the GPU parallel implementation of Ehrlich-Aberth method.

 
 \begin{enumerate}
 \begin{algorithm}[htpb]
 
 \begin{enumerate}
 \begin{algorithm}[htpb]


@@ -657,10 +655,9 @@ Algorithm~\ref{alg2-cuda} shows a sketch of the Ehrlich-Aberth method using CUDA
  
 \section{The EA algorithm on Multi-GPU}
 
  
 \section{The EA algorithm on Multi-GPU}
 

-\subsection{MGPU (OpenMP-CUDA) approach}
-Our OpenMP-CUDA implementation of EA algorithm is based on the hybrid OpenMP and CUDA programming model. It works
-as follows.
-Based on the metadata, a shared memory is used to make data evenly shared among OpenMP threads. The shared data are the solution vector $Z$, the polynomial to solve $P$. vector of error of stop condition $\Delta z$. Let(T\_omp) number of OpenMP threads is equal to the number of GPUs, each threads OpenMP checks one GPU,  and control a part of the shared memory, that is a part of the vector Z  like: $(n/num\_gpu)$ roots, n: the polynomial's degrees, $num\_gpu$ the number of GPUs. Each OpenMP thread copies its data from host memory to GPU’s device memory.Than every GPU will have a grid of computation organized with its performances and the size of data of which it checks and compute kernels. %In principle a grid is set by two parameter DimGrid, the number of block per grid, DimBloc: the number of threads per block. The following schema  shows the architecture of (CUDA,OpenMP).
+\subsection{MGPU : an OpenMP-CUDA approach}
+Our OpenMP-CUDA implementation of EA algorithm is based on the hybrid OpenMP and CUDA programming model. It works as follows.
+Based on the metadata, a shared memory is used to make data evenly shared among OpenMP threads. The shared data are the solution vector $Z$, the polynomial to solve $P$, and  the error vector $\Delta z$. Let (T\_omp) the number of OpenMP threads be equal to the number of GPUs, each OpenMP thread binds to one GPU,  and controls a part of the shared memory, that is a part of the vector Z , that is $(n/num\_gpu)$ roots where $n$ is the polynomial's degree and $num\_gpu$ the total number of available GPUs. Each OpenMP thread copies its data from host memory to GPU’s device memory.Then every GPU will have a grid of computation organized according to the device performance and the size of data on which it runs the computation kernels. %In principle a grid is set by two parameter DimGrid, the number of block per grid, DimBloc: the number of threads per block. The following schema  shows the architecture of (CUDA,OpenMP).

 
 %\begin{figure}[htbp]
 %\centering
 
 %\begin{figure}[htbp]
 %\centering


@@ -670,7 +667,7 @@ Based on the metadata, a shared memory is used to make data evenly shared among
 %\end{figure}
 %Each thread OpenMP compute the kernels on GPUs,than after each iteration they copy out the data from GPU memory to CPU shared memory. The kernels are re-runs is up to the roots converge sufficiently. Here are below the corresponding algorithm:
 
 %\end{figure}
 %Each thread OpenMP compute the kernels on GPUs,than after each iteration they copy out the data from GPU memory to CPU shared memory. The kernels are re-runs is up to the roots converge sufficiently. Here are below the corresponding algorithm:
 

-$num\_gpus$ thread OpenMP are created using \verb=omp_set_num_threads();=function (line,Algorithm \ref{alg2-cuda-openmp}), the shared memory is created using \verb=#pragma omp parallel shared()= OpenMP function (line 5,Algorithm\ref{alg2-cuda-openmp}), than each OpenMP threads allocate and copy initial data from CPU memory to the GPU global memories, execute the kernels on GPU, and compute only his portion of roots indicated with variable \textit{index} initialized in (line 5, Algorithm \ref{alg2-cuda-openmp}), used as input data in the $kernel\_update$ (line 10, Algorithm \ref{alg2-cuda-openmp}). After each iteration, OpenMP threads synchronize using \verb=#pragma omp barrier;= to recuperate all values of vector $\Delta z$, to compute the maximum stop condition in vector $\Delta z$(line 12, Algorithm \ref{alg2-cuda-openmp}).Finally,they copy the results from GPU memories to CPU memory. The OpenMP threads execute kernels until the roots converge sufficiently.  
+$num\_gpus$ OpenMP threads  are created using \verb=omp_set_num_threads();=function (step $3$, Algorithm \ref{alg2-cuda-openmp}), the shared memory is created using \verb=#pragma omp parallel shared()= OpenMP function (line $5$, Algorithm\ref{alg2-cuda-openmp}), then each OpenMP thread allocates memory and copies initial data from CPU memory to GPU global memory, executes the kernels on GPU, but computes only his portion of roots indicated with variable \textit{index} initialized in (line 5, Algorithm \ref{alg2-cuda-openmp}), used as input data in the $kernel\_update$ (line 10, Algorithm \ref{alg2-cuda-openmp}). After each iteration, all OpenMP threads synchronize using \verb=#pragma omp barrier;= to gather all the correct values of $\Delta z$, thus allowing the computation the maximum stop condition on vector $\Delta z$ (line 12, Algorithm \ref{alg2-cuda-openmp}). Finally, threads copy the results from GPU memories to CPU memory. The OpenMP threads execute kernels until the roots sufficiently converge.  

 \begin{enumerate}
 \begin{algorithm}[htpb]
 \label{alg2-cuda-openmp}
 \begin{enumerate}
 \begin{algorithm}[htpb]
 \label{alg2-cuda-openmp}


@@ -678,14 +675,14 @@ $num\_gpus$ thread OpenMP are created using \verb=omp_set_num_threads();=functio
 \caption{CUDA-OpenMP Algorithm to find roots with the Ehrlich-Aberth method}
 
 \KwIn{$Z^{0}$ (Initial root's vector), $\varepsilon$ (Error tolerance
 \caption{CUDA-OpenMP Algorithm to find roots with the Ehrlich-Aberth method}
 
 \KwIn{$Z^{0}$ (Initial root's vector), $\varepsilon$ (Error tolerance

-  threshold), P (Polynomial to solve), Pu (Derivative of P), $n$ (Polynomial degrees), $\Delta z$ ( Vector of errors of stop condition), $num_gpus$ (number of OpenMP threads/ number of GPUs), $Size$ (number of roots)}
+  threshold), P (Polynomial to solve), Pu (Derivative of P), $n$ (Polynomial degree), $\Delta z$ ( Vector of errors for stop condition), $num_gpus$ (number of OpenMP threads/ Number of GPUs), $Size$ (number of roots)}

 
 

-\KwOut {$Z$ (Solution root's vector), $ZPrec$ (Previous solution root's vector)}
+\KwOut {$Z$ ( Root's vector), $ZPrec$ (Previous root's vector)}

 
 \BlankLine
 
 
 \BlankLine
 

-\item Initialization of the of P\;
-\item Initialization of the of Pu\;
+\item Initialization of P\;
+\item Initialization of Pu\;

 \item Initialization of the solution vector $Z^{0}$\;
 \verb=omp_set_num_threads(num_gpus);=
 \verb=#pragma omp parallel shared(Z,$\Delta$ z,P);=
 \item Initialization of the solution vector $Z^{0}$\;
 \verb=omp_set_num_threads(num_gpus);=
 \verb=#pragma omp parallel shared(Z,$\Delta$ z,P);=


@@ -710,16 +707,16 @@ $num\_gpus$ thread OpenMP are created using \verb=omp_set_num_threads();=functio
 
 
 
 
 
 

-\subsection{Multi-GPU (MPI-CUDA) approach}
+\subsection{Multi-GPU : an MPI-CUDA approach}

 %\begin{figure}[htbp]
 %\centering
  % \includegraphics[angle=-90,width=0.2\textwidth]{MPI-CUDA}
 %\caption{The MPI-CUDA architecture }
 %\label{fig:03}
 %\end{figure}
 %\begin{figure}[htbp]
 %\centering
  % \includegraphics[angle=-90,width=0.2\textwidth]{MPI-CUDA}
 %\caption{The MPI-CUDA architecture }
 %\label{fig:03}
 %\end{figure}

-Our parallel implementation of the Ehrlich-Aberth method to find root polynomial using (CUDA-MPI) approach, splits input data of the polynomial to solve between MPI processes. From Algorithm 3, the input data are the polynomial to solve $P$, the solution vector $Z$, the previous solution vector $zPrev$, and the Value of errors of stop condition $\Delta z$. Let $p$ denote the number of MPI processes on and $n$ the size of the polynomial to be solved. The algorithm performs a simple data partitioning by creating $p$ portions, of at most $⌈n/p⌉$ roots to find per MPI process, for each element mentioned above. Consequently, each MPI process $k$ will have its own solution vector $Z_{k}$,polynomial to be solved $p_{k}$, the error of stop condition $\Delta z_{k}$, Than each MPI processes compute only $⌈n/p⌉$ roots.
+Our parallel implementation of the Ehrlich-Aberth method to find root of polynomials using a CUDA-MPI approach, splits input data of the polynomial to solve among MPI processes. In Algorithm \ref{alg2-cuda-mpi}, input data are the polynomial to solve $P$, the solution vector $Z$, the previous solution vector $ZPrev$, and the value of errors of stop condition $\Delta z$. Let $p$ denote the number of MPI processes on and $n$ the degree of the polynomial to be solved. The algorithm performs a simple data partitioning by creating $p$ portions, of at most $⌈n/p⌉$ roots to find per MPI process, for each $Z$ and $ZPrec$. Consequently, each MPI process of rank $k$ will have its own solution vector $Z_{k}$ and $ZPrec$, the error related to the stop condition $\Delta z_{k}$, enabling each MPI process to compute $⌈n/p⌉$ roots.

 
 

-Since a GPU works only on data of its memory, all local input data, $Z_{k}, p_{k}$ and $\Delta z_{k}$, must be transferred from CPU memories to the corresponding GPU memories. Afterward, the same EA algorithm (Algorithm 1) is run by all processes but on different sub-polynomial root $ p(x)_{k}=\sum_{i=k(\frac{n}{p})}^{k+1(\frac{n}{p})} a_{i}x^{i}, k=1,...,p$.  Each processes MPI execute the  loop \verb=(While(...)...do)= contain the kernels. Than each process MPI  compute only his portion of roots indicated with variable \textit{index} initialized in (line 5, Algorithm \ref{alg2-cuda-mpi}), used as input data in the $kernel\_update$ (line 10, Algorithm \ref{alg2-cuda-mpi}). After each iteration, MPI processes synchronize using \verb=MPI_Allreduce= function, in order to compute the maximum error stops condition $\Delta z_{k}$ computed by each process MPI line (line, Algorithm\ref{alg2-cuda-mpi}), and  copy the values of new roots computed from GPU memories to CPU memories, than communicate  her results to the neighboring processes,using \verb=MPI_Alltoallv=. If maximum stop condition $error > \epsilon$ the processes stay to execute the loop \verb= while(...)...do= until all the roots converge sufficiently.
+Since a GPU works only on data already allocated in its memory, all local input data, $Z_{k}$, $ZPrec$ and $\Delta z_{k}$, must be transferred from CPU memories to the corresponding GPU memories. Afterwards, the same EA algorithm (Algorithm \ref{alg1-cuda}) is run by all processes but on different sub-polynomial roots $ p(x)_{k}=\sum_{i=1}^{n} a_{i}x^{i}, k=1,...,p$.  Each MPI process executes the  loop \verb=(While(...)...do)= containing the CUDA kernels. Then each MPI process  computes only its own portion of the roots indicated with variable \textit{index} initialized in (line 5, Algorithm \ref{alg2-cuda-mpi}), used as an input variable in the $kernel\_update$ (line 10, Algorithm \ref{alg2-cuda-mpi}). After each iteration, MPI processes synchronize using \verb=MPI_Allreduce= function, in order to compute the maximum error related to the stop condition; the reduction on $\Delta z_{k}$ by each MPI process on (line, Algorithm\ref{alg2-cuda-mpi}), and  copy the values of new computed roots  from GPU memories to CPU memories, then communicate their results to other processes,using \verb=MPI_Alltoall=. If the stop condition is not verified ($error > \epsilon$) then processes stay withing the loop \verb= while(...)...do= until all the roots sufficiently converge.

 
 \begin{enumerate}
 \begin{algorithm}[htpb]
 
 \begin{enumerate}
 \begin{algorithm}[htpb]


@@ -733,24 +730,22 @@ Since a GPU works only on data of its memory, all local input data, $Z_{k}, p_{k
 \KwOut {$Z$ (Solution root's vector), $ZPrec$ (Previous solution root's vector)}
 
 \BlankLine
 \KwOut {$Z$ (Solution root's vector), $ZPrec$ (Previous solution root's vector)}
 
 \BlankLine

-\item Initialization of the P\;
-\item Initialization of the Pu\;
+\item Initialization of P\;
+\item Initialization of Pu\;

 \item Initialization of the solution vector $Z^{0}$\;
 \item Initialization of the solution vector $Z^{0}$\;

-\item Allocate and copy initial data from CPU memories to the GPU global memories\;
+\item Allocate and copy initial data from CPU memories to GPU global memories\;

 \item $index= Size/num_gpus$\;
 \item k=0\;
 \While {$error > \epsilon$}{
 \item Let $\Delta z=0$\;
 \item $index= Size/num_gpus$\;
 \item k=0\;
 \While {$error > \epsilon$}{
 \item Let $\Delta z=0$\;

-\item $ kernel\_save(ZPrec,Z)$\;
+\item $kernel\_save(ZPrec,Z)$\;

 \item  k=k+1\;
 \item  k=k+1\;

-\item $ kernel\_update(Z,P,Pu,index)$\;
+\item $kernel\_update(Z,P,Pu,index)$\;

 \item $kernel\_testConverge(\Delta z,Z,ZPrec)$\;
 \item ComputeMaxError($\Delta z$,error)\;
 \item Copy results from GPU memories to CPU memories\;
 \item $kernel\_testConverge(\Delta z,Z,ZPrec)$\;
 \item ComputeMaxError($\Delta z$,error)\;
 \item Copy results from GPU memories to CPU memories\;

-\item Send $Z[id]$ to all neighboring processes\;
-\item Receive $Z[j]$ from neighboring process j\;
-
-
+\item Send $Z[id]$ to all processes\;
+\item Receive $Z[j]$ from every other process j\;

 }
 \end{algorithm}
 \end{enumerate}
 }
 \end{algorithm}
 \end{enumerate}


@@ -964,12 +959,13 @@ This is due to the use of parallelization MPI paradigm that divides the polynomi
 
 
 \section{Conclusion}
 
 
 \section{Conclusion}

-In this paper, we have presented a parallel implementation of Ehrlich-Aberth algorithm for solving full and sparse polynomials, on single GPU with CUDA and Multi-GPUs using two parallel paradigm, shared memory with OpenMP, distributed memory with MPI.(CUDA-OpenMP) approach and (CUDA-MPI) approach, 
-We have performed many experiments with the Ehrlich-Aberth method in single GPU, Multi-GPU with (CUDA-OpenMP) approach, Multi-GPU with (CUDA-MPI) approach for sparse and full polynomials. the experiments show that, using parallel programming model like (OpenMP, MPI) can effectively manage multiple graphics cards to work together to solve the same problem and accelerate parallel applications, like (CUDA MPI) approach with 4 GPUs can solve a polynomial of 1,000,000 4 speed up than on single GPU.
+In this paper, we have presented a parallel implementation of Ehrlich-Aberth algorithm for solving full and sparse polynomials, on single GPU with CUDA and on Multi-GPUs using two parallel paradigm, shared memory with OpenMP, distributed memory with MPI.(CUDA-OpenMP) approach and (CUDA-MPI) approach, 
+We have performed many experiments with the Ehrlich-Aberth method in single GPU, Multi-GPU with (CUDA-OpenMP) approach, Multi-GPU with (CUDA-MPI) approach for sparse and full polynomials. the experiments show that, using parallel programming model like (OpenMP, MPI) can efficiently manage multiple graphics cards to work together to solve the same problem and accelerate parallel applications, like (CUDA MPI) approach with 4 GPUs can solve a polynomial of 1,000,000 4 speed up than on single GPU.

 
 
 
 

-In future, we will evaluate our parallel implementation of Ehrlich-Aberth algorithm on other parallel programming model 
+%In future, we will evaluate our parallel implementation of Ehrlich-Aberth algorithm on other parallel programming model 

 
 

+Our next objective is to extend the model presented here at nodes clusters frame multi-GPU , with a three-level scheme: inter-node communication via MPI processes (distributed memory), management of multi-GPU node by OpenMP threads (shared memory).

 
 %present a communication approach between multiple GPUs. The comparison between MPI and OpenMP as GPUs controllers shows that these
 %solutions can effectively manage multiple graphics cards to work together
 
 %present a communication approach between multiple GPUs. The comparison between MPI and OpenMP as GPUs controllers shows that these
 %solutions can effectively manage multiple graphics cards to work together