% *** GRAPHICS RELATED PACKAGES ***
%
\ifCLASSINFOpdf
- % \usepackage[pdftex]{graphicx}
+ \usepackage[pdftex]{graphicx}
+
% declare the path(s) where your graphic files are
% \graphicspath{{../pdf/}{../jpeg/}}
% and their extensions so you won't have to specify these with
% *** MATH PACKAGES ***
%
-%\usepackage{amsmath}
+\usepackage{amsmath}
% A popular package from the American Mathematical Society that provides
% many useful and powerful commands for dealing with mathematics.
%
% *** SPECIALIZED LIST PACKAGES ***
%
-%\usepackage{algorithmic}
+\usepackage{algorithmic}
% algorithmic.sty was written by Peter Williams and Rogerio Brito.
% This package provides an algorithmic environment fo describing algorithms.
% You can use the algorithmic environment in-text or within a figure
% Also of interest may be the (relatively newer and more customizable)
% algorithmicx.sty package by Szasz Janos:
% http://www.ctan.org/pkg/algorithmicx
-
+\usepackage[ruled,vlined]{algorithm2e}
% correct bad hyphenation here
\hyphenation{op-tical net-works semi-conduc-tor}
-\usepackage{graphicx}
+%\usepackage{graphicx}
+
\begin{document}
%
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 manage CUDA context of different GPUs. A direct method for controlling the various GPU is to use as many threads or processes that GPU. We can choose the GPU index based on the identifier of OpenMP thread or the rank of the MPI process. Both approaches will be created.
\section{The EA algorithm on single GPU}
+\subsection{the EA method}
+%\begin{figure}[htbp]
+%\centering
+ % \includegraphics[angle=-90,width=0.5\textwidth]{EA-Algorithm}
+%\caption{The Ehrlich-Aberth algorithm on single GPU}
+%\label{fig:03}
+%\end{figure}
+
+the Ehrlich-Aberth method is an iterative method, contain 4 steps, start from the initial approximations of all the
+roots of the polynomial,the second step initialize the solution vector $Z$ using the Guggenheimer method to assure the distinction of the initial vector roots, than in step 3 we apply the the iterative function based on the Newton's method and Weiestrass operator[...,...], wich will make it possible to converge to the roots solution, provided that all the root are different. At the end of each application of the iterative function, a stop condition is verified consists in stopping the iterative process when the whole of the modules of the roots
+are lower than a fixed value $ε$
+\subsection{EA parallel implementation on CUDA}
+Like any parallel code, a GPU parallel implementation first
+requires to determine the sequential tasks and the
+parallelizable parts of the sequential version of the
+program/algorithm. In our case, all the operations that are easy
+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
+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
+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
+be executed on the GPU for all the operations dealing with the
+computation on complex numbers and the evaluation of the
+polynomials. As said previously, we managed both functions
+of evaluation of a polynomial: the normal method, based on
+the method of Horner and the method based on the logarithm
+of the polynomial. All these methods were rather long to
+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
+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.
+
+\begin{enumerate}
+\begin{algorithm}[htpb]
+\label{alg2-cuda}
+%\LinesNumbered
+\caption{CUDA 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_{max}$ (Maximum value of stop condition)}
+
+\KwOut {$Z$ (Solution root's vector), $ZPrec$ (Previous solution root's vector)}
+
+\BlankLine
+
+\item Initialization of the of P\;
+\item Initialization of the of Pu\;
+\item Initialization of the solution vector $Z^{0}$\;
+\item Allocate and copy initial data to the GPU global memory\;
+\item k=0\;
+\While {$\Delta z_{max} > \epsilon$}{
+\item Let $\Delta z_{max}=0$\;
+\item $ kernel\_save(ZPrec,Z)$\;
+\item k=k+1\;
+\item $ kernel\_update(Z,P,Pu)$\;
+\item $kernel\_testConverge(\Delta z_{max},Z,ZPrec)$\;
+
+}
+\item Copy results from GPU memory to CPU memory\;
+\end{algorithm}
+\end{enumerate}
+~\\
+
\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).
+
+%\begin{figure}[htbp]
+%\centering
+ % \includegraphics[angle=-90,width=0.5\textwidth]{OpenMP-CUDA}
+%\caption{The OpenMP-CUDA architecture}
+%\label{fig:03}
+%\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.
+\begin{enumerate}
+\begin{algorithm}[htpb]
+\label{alg2-cuda-openmp}
+%\LinesNumbered
+\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)}
+
+\KwOut {$Z$ (Solution root's vector), $ZPrec$ (Previous solution root's vector)}
+
+\BlankLine
+
+\item Initialization of the of P\;
+\item Initialization of the 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);=
+\verb=cudaGetDevice(gpu_id);=
+\item Allocate and copy initial data from CPU memory to the GPU global memories\;
+\item index= $Size/num\_gpus$\;
+\item k=0\;
+\While {$error > \epsilon$}{
+\item Let $\Delta z=0$\;
+\item $ kernel\_save(ZPrec,Z)$\;
+\item k=k+1\;
+\item $ kernel\_update(Z,P,Pu,index)$\;
+\item $kernel\_testConverge(\Delta z[gpu\_id],Z,ZPrec)$\;
+%\verb=#pragma omp barrier;=
+\item error= Max($\Delta z$)\;
+}
+
+\item Copy results from GPU memories to CPU memory\;
+\end{algorithm}
+\end{enumerate}
+~\\
+
+
+
+\subsection{Multi-GPU (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}
+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.
+
+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.
+
+\begin{enumerate}
+\begin{algorithm}[htpb]
+\label{alg2-cuda-mpi}
+%\LinesNumbered
+\caption{CUDA-MPI 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$ ( error of stop condition), $num_gpus$ (number of MPI processes/ number of GPUs), Size (number of roots)}
+
+\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 the solution vector $Z^{0}$\;
+\item Allocate and copy initial data from CPU memories to the GPU global memories\;
+\item $index= Size/num_gpus$\;
+\item k=0\;
+\While {$error > \epsilon$}{
+\item Let $\Delta z=0$\;
+\item $ kernel\_save(ZPrec,Z)$\;
+\item k=k+1\;
+\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 Send $Z[id]$ to all neighboring processes\;
+\item Receive $Z[j]$ from neighboring process j\;
+
+
+}
+\end{algorithm}
+\end{enumerate}
+~\\
-\subsection{MGPU (OpenMP-CUDA)approach}
-\subsection{MGPU (MPI-CUDA)approach}
\section{experiments}
+We study two categories of polynomials: sparse polynomials and full polynomials.\\
+{\it A sparse polynomial} is a polynomial for which only some coefficients are not null. In this paper, we consider sparse polynomials for which the roots are distributed on 2 distinct circles:
+\begin{equation}
+ \forall \alpha_{1} \alpha_{2} \in C,\forall n_{1},n_{2} \in N^{*}; P(z)= (z^{n_{1}}-\alpha_{1})(z^{n_{2}}-\alpha_{2})
+\end{equation}\noindent
+{\it A full polynomial} is, in contrast, a polynomial for which all the coefficients are not null. A full polynomial is defined by:
+%%\begin{equation}
+ %%\forall \alpha_{i} \in C,\forall n_{i}\in N^{*}; P(z)= \sum^{n}_{i=1}(z^{n^{i}}.a_{i})
+%%\end{equation}
+
+\begin{equation}
+ {\Large \forall a_{i} \in C, i\in N; p(x)=\sum^{n}_{i=0} a_{i}.x^{i}}
+\end{equation}
+For our tests, a CPU Intel(R) Xeon(R) CPU E5620@2.40GHz and a GPU K40 (with 6 Go of ram) are used.
+
+We performed a set of experiments on single GPU and Multi-GPU using (OpenMP/MPI) to find roots polynomials with EA algorithm, for both sparse and full polynomials of different sizes. We took into account the execution times and the polynomial size performed by sum or each experiment.
+All experimental results obtained from the simulations are made in
+double precision data, the convergence threshold of the methods is set
+to $10^{-7}$.
+%Since we were more interested in the comparison of the
+%performance behaviors of Ehrlich-Aberth and Durand-Kerner methods on
+%CPUs versus on GPUs.
+The initialization values of the vector solution
+of the methods are given in %Section~\ref{sec:vec_initialization}.
\begin{figure}[htbp]
\centering
- \includegraphics[angle=-90,width=0.8\textwidth]{GPU_openmp}
-\caption{Execution times in seconds of the Ehrlich-Aberth method on GPUs using shared memory paradigm with OpenMP}
+ \includegraphics[angle=-90,width=0.5\textwidth]{Sparse_openmp}
+\caption{Execution times in seconds of the Ehrlich-Aberth method for solving sparse polynomials on GPUs using shared memory paradigm with OpenMP}
\label{fig:01}
\end{figure}
\begin{figure}[htbp]
\centering
- \includegraphics[angle=-90,width=0.8\textwidth]{GPU_mpi}
-\caption{Execution times in seconds of the Ehrlich-Aberth method on GPUs using distributed memory paradigm with MPI}
+ \includegraphics[angle=-90,width=0.5\textwidth]{Sparse_mpi}
+\caption{Execution times in seconds of the Ehrlich-Aberth method for solving sparse polynomials on GPUs using distributed memory paradigm with MPI}
\label{fig:02}
\end{figure}
+
+\begin{figure}[htbp]
+\centering
+ \includegraphics[angle=-90,width=0.5\textwidth]{Full_openmp}
+\caption{Execution times in seconds of the Ehrlich-Aberth method for solving full polynomials on GPUs using shared memory paradigm with OpenMP}
+\label{fig:03}
+\end{figure}
+
+\begin{figure}[htbp]
+\centering
+ \includegraphics[angle=-90,width=0.5\textwidth]{Full_mpi}
+\caption{Execution times in seconds of the Ehrlich-Aberth method for full polynomials on GPUs using distributed memory paradigm with MPI}
+\label{fig:04}
+\end{figure}
+
+\begin{figure}[htbp]
+\centering
+ \includegraphics[angle=-90,width=0.5\textwidth]{Sparse_mpivsomp}
+\caption{Comparaison between MPI and OpenMP versions of the Ehrlich-Aberth method for solving sparse plynomials on GPUs}
+\label{fig:05}
+\end{figure}
+
+\begin{figure}[htbp]
+\centering
+ \includegraphics[angle=-90,width=0.5\textwidth]{Full_mpivsomp}
+\caption{Comparaison between MPI and OpenMP versions of the Ehrlich-Aberth method for solving full polynomials on GPUs}
+\label{fig:06}
+\end{figure}
+
+\begin{figure}[htbp]
+\centering
+ \includegraphics[angle=-90,width=0.5\textwidth]{MPI_mpivsomp}
+\caption{Comparaison of execution times of the Ehrlich-Aberth method for solving sparse and full polynomials on GPUs with distributed memory paradigm using MPI}
+\label{fig:07}
+\end{figure}
+
+\begin{figure}[htbp]
+\centering
+ \includegraphics[angle=-90,width=0.5\textwidth]{OMP_mpivsomp}
+\caption{Comparaison of execution times of the Ehrlich-Aberth method for solving sparse and full polynomials on GPUs with shared memory paradigm using OpenMP}
+\label{fig:08}
+\end{figure}
+
% An example of a floating figure using the graphicx package.
% Note that \label must occur AFTER (or within) \caption.
% For figures, \caption should occur after the \includegraphics.
