\section{The EA algorithm on Multi-GPU}
-\subsection{MGPU (OpenMP-CUDA)approach}
-Before starting computations, our parallel implementation shared input data of the root finding polynomial between OpenMP threads. From Algorithm 1, the input data are the solution vector $Z$, the polynomial to solve $P$. Let number of OpenMP threads is equal to the number of GPUs, each threads OpenMP ( T-omp) checks one GPU, and control a part of the shared memory, that is a part of the vector Z like: $(n/Nbr_gpu)$ roots, n: the polynomial's degrees, $Nbr_gpu$ the number of GPUs. Then every GPU will have a grid of computation organized with its performances and the size of data of which it checks. 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 (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
%\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:
-
-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}
+\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)}
+ 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)}
\item Initialization of the of Pu\;
\item Initialization of the solution vector $Z^{0}$\;
\verb=omp_set_num_threads(num_gpus);=
-\verb=cudaGetDevice(gpu_id);=
\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\;
-\subsection{Multi-GPU (MPI-CUDA)approach}
+\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}
+\label{alg2-cuda-mpi}
%\LinesNumbered
\caption{CUDA-MPI Algorithm to find roots with the Ehrlich-Aberth method}
\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\;
-\item ComputeMaxError($\Delta z$,error)\;
+
}
\end{algorithm}
~\\
\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.5\textwidth]{Sparse_openmp}
