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
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}.
\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:
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
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]
\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
%\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}
\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
-\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);=
-\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}
-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]
\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 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 $ kernel\_save(ZPrec,Z)$\;
+\item $kernel\_save(ZPrec,Z)$\;
\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 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}
