new

[book_gpu.git] / BookGPU / Chapters / chapter13 / ch13.tex

diff --git a/BookGPU/Chapters/chapter13/ch13.tex b/BookGPU/Chapters/chapter13/ch13.tex
index 89e2ce73adfc7e59d93c63cc4cf6bb510e632826..ce2c7c281c5162b463f25a3dea8c0d23760a2937 100755 (executable)
--- a/BookGPU/Chapters/chapter13/ch13.tex
+++ b/BookGPU/Chapters/chapter13/ch13.tex
@@ -299,12 +299,12 @@ and $\forall j\in\{1,\ldots,\alpha\}$,
 \label{ch13:eq:15}
 \end{equation}
 
 \label{ch13:eq:15}
 \end{equation}
 

-The previous asynchronous scheme\index{asynchronous} of the projected Richardson
+The previous asynchronous scheme\index{asynchronous iterations} of the projected Richardson

 method models computations that are carried out in parallel without order or
 synchronization (according to the behavior of the parallel iterative method) and
 describes a subdomain method without overlapping. It is a general model that takes
 into account all possible situations of parallel computations and nonblocking message
 method models computations that are carried out in parallel without order or
 synchronization (according to the behavior of the parallel iterative method) and
 describes a subdomain method without overlapping. It is a general model that takes
 into account all possible situations of parallel computations and nonblocking message

-passing. So, the synchronous iterative scheme\index{synchronous} is defined by
+passing. So, the synchronous iterative scheme\index{synchronous iterations} is defined by

 \begin{equation}
 \forall j\in\{1,\ldots,\alpha\} \mbox{,~} \forall p\in\mathbb{N} \mbox{,~} \rho_j(p)=p.
 \label{ch13:eq:16}
 \begin{equation}
 \forall j\in\{1,\ldots,\alpha\} \mbox{,~} \forall p\in\mathbb{N} \mbox{,~} \rho_j(p)=p.
 \label{ch13:eq:16}


@@ -322,7 +322,7 @@ each component of the vector must be relaxed an infinite number of times. The ch
 relaxed components to be used in the computational process may be guided by any criterion,
 and in particular, a natural criterion is to pickup the most recently available
 values of the components computed by the processors. Furthermore, the asynchronous
 relaxed components to be used in the computational process may be guided by any criterion,
 and in particular, a natural criterion is to pickup the most recently available
 values of the components computed by the processors. Furthermore, the asynchronous

-iterations are implemented by means of nonblocking MPI communication subroutines\index{MPI subroutines!nonblocking}
+iterations are implemented by means of nonblocking MPI communication subroutines\index{MPI!nonblocking}

 (asynchronous communications).
 
 The important property ensuring the convergence of the parallel projected Richardson
 (asynchronous communications).
 
 The important property ensuring the convergence of the parallel projected Richardson


@@ -477,7 +477,7 @@ The first operation of this function is implemented as kernels to be performed b
 \end{itemize}
 As mentioned previously, we develop the \emph{synchronous} and \emph{asynchronous}
 algorithms of the projected Richardson method. Obviously, in this scope, the
 \end{itemize}
 As mentioned previously, we develop the \emph{synchronous} and \emph{asynchronous}
 algorithms of the projected Richardson method. Obviously, in this scope, the

-synchronous\index{synchronous} or asynchronous\index{asynchronous} communications
+synchronous\index{synchronous iterations} or asynchronous\index{asynchronous iterations} communications

 refer to the communications between the CPU cores (MPI processes) on the GPU cluster,
 in order to exchange the vector elements associated to subdomain boundaries. For
 the memory copies between a CPU core and its GPU, we use the synchronous communication
 refer to the communications between the CPU cores (MPI processes) on the GPU cluster,
 in order to exchange the vector elements associated to subdomain boundaries. For
 the memory copies between a CPU core and its GPU, we use the synchronous communication


@@ -486,13 +486,13 @@ in the synchronous algorithm and the asynchronous ones: \verb+cublasSetVectorAsy
 and \verb+cublasGetVectorAsync()+ in the asynchronous algorithm. Moreover, we
 use the communication routines of the MPI library to carry out the data exchanges
 between the neighboring nodes. We use the following communication routines: \verb+MPI_Isend()+
 and \verb+cublasGetVectorAsync()+ in the asynchronous algorithm. Moreover, we
 use the communication routines of the MPI library to carry out the data exchanges
 between the neighboring nodes. We use the following communication routines: \verb+MPI_Isend()+

-and \verb+MPI_Irecv()+ to perform nonblocking\index{MPI subroutines!nonblocking}
+and \verb+MPI_Irecv()+ to perform nonblocking\index{MPI!nonblocking}

 sends and receives, respectively. For the synchronous algorithm, we use the MPI
 routine \verb+MPI_Waitall()+ which puts the MPI process of a computing node in
 blocking status until all data exchanges with neighboring nodes (sends and receives)
 are completed. In contrast, for the asynchronous algorithms, we use the MPI routine
 \verb+MPI_Test()+ which tests the completion of a data exchange (send or receives)
 sends and receives, respectively. For the synchronous algorithm, we use the MPI
 routine \verb+MPI_Waitall()+ which puts the MPI process of a computing node in
 blocking status until all data exchanges with neighboring nodes (sends and receives)
 are completed. In contrast, for the asynchronous algorithms, we use the MPI routine
 \verb+MPI_Test()+ which tests the completion of a data exchange (send or receives)

-without putting the MPI process in blocking status\index{MPI subroutines!blocking}.   
+without putting the MPI process in blocking status\index{MPI!blocking}.   

 
 The function $Compute\_New\_Vector\_Elements()$ (line~$6$ in Algorithm~\ref{ch13:alg:02})
 computes, at each iteration, the new elements of the iterate vector $U$. Its general code
 
 The function $Compute\_New\_Vector\_Elements()$ (line~$6$ in Algorithm~\ref{ch13:alg:02})
 computes, at each iteration, the new elements of the iterate vector $U$. Its general code


@@ -587,7 +587,7 @@ of relaxations is possible in both synchronous and asynchronous cases. On the
 other hand, an iteration is the update of at least all vector components with
 $F_i$.
 
 other hand, an iteration is the update of at least all vector components with
 $F_i$.
 

-In the synchronous\index{synchronous} algorithm, the global convergence is detected
+In the synchronous\index{synchronous iterations} algorithm, the global convergence is detected

 when the maximal value of the absolute error, $error$, is sufficiently small and/or
 the maximum number of relaxations, $MaxRelax$, is reached, as follows:
 $$
 when the maximal value of the absolute error, $error$, is sufficiently small and/or
 the maximum number of relaxations, $MaxRelax$, is reached, as follows:
 $$


@@ -598,11 +598,11 @@ AllReduce(error,\hspace{0.1cm}maxerror,\hspace{0.1cm}MAX); \\
 conv \leftarrow true;
 \end{array}
 $$
 conv \leftarrow true;
 \end{array}
 $$

-where the function $AllReduce()$ uses the MPI global reduction subroutine\index{MPI subroutines!global}
+where the function $AllReduce()$ uses the MPI global reduction subroutine\index{MPI!global}

 \verb+MPI_Allreduce()+ to compute the maximal value, $maxerror$, among the local
 absolute errors, $error$, of all computing nodes, and $p$ (in Algorithm~\ref{ch13:alg:02})
 is used as a counter of the local relaxations carried out by a computing node. In
 \verb+MPI_Allreduce()+ to compute the maximal value, $maxerror$, among the local
 absolute errors, $error$, of all computing nodes, and $p$ (in Algorithm~\ref{ch13:alg:02})
 is used as a counter of the local relaxations carried out by a computing node. In

-the asynchronous\index{asynchronous} algorithms, the global convergence is detected
+the asynchronous\index{asynchronous iterations} algorithms, the global convergence is detected

 when all computing nodes locally converge. For this, we use a token ring architecture
 around which a boolean token travels, in one direction, from a computing node to another.
 Starting from node $0$, the boolean token is set to $true$ by node $i$ if the local
 when all computing nodes locally converge. For this, we use a token ring architecture
 around which a boolean token travels, in one direction, from a computing node to another.
 Starting from node $0$, the boolean token is set to $true$ by node $i$ if the local


@@ -698,6 +698,7 @@ $800^{3}$                     & $222,108.09$       & $1,769,232$      & $188,790
 
 \begin{table}
 \centering
 
 \begin{table}
 \centering

+\begin{scriptsize}

 \begin{tabular}{|c|c|c|c|c|c|c|c|}
 \hline
 \multirow{2}{*}{\bf Pb. size} & \multicolumn{3}{c|}{\bf Synchronous}                 & \multicolumn{3}{c|}{\bf Asynchronous}                & \multirow{2}{*}{\bf Gain\%}  \\ \cline{2-7}
 \begin{tabular}{|c|c|c|c|c|c|c|c|}
 \hline
 \multirow{2}{*}{\bf Pb. size} & \multicolumn{3}{c|}{\bf Synchronous}                 & \multicolumn{3}{c|}{\bf Asynchronous}                & \multirow{2}{*}{\bf Gain\%}  \\ \cline{2-7}


@@ -712,6 +713,7 @@ $768^{3}$                    & $4,112.68$         & $831,144$        & $50.13$
 
 $800^{3}$                    & $3,950.87$         & $899,088$        & $56.22$         & $3,636.57$        & $834,900$        & $51.91$     & $7.95$ \\ \hline 
 \end{tabular}
 
 $800^{3}$                    & $3,950.87$         & $899,088$        & $56.22$         & $3,636.57$        & $834,900$        & $51.91$     & $7.95$ \\ \hline 
 \end{tabular}

+\end{scriptsize}

 \vspace{0.5cm}
 \caption{Execution times in seconds of the parallel projected Richardson method implemented on a cluster of 12 GPUs.}
 \label{ch13:tab:02}
 \vspace{0.5cm}
 \caption{Execution times in seconds of the parallel projected Richardson method implemented on a cluster of 12 GPUs.}
 \label{ch13:tab:02}


@@ -745,7 +747,7 @@ consequently it also depends on the number of computing nodes.
 %%--------------------------%%
 \section{Red-black ordering technique}
 \label{ch13:sec:06}
 %%--------------------------%%
 \section{Red-black ordering technique}
 \label{ch13:sec:06}

-As is wellknown, the Jacobi method\index{iterative method!Jacobi} is characterized
+As is well-known, the Jacobi method\index{iterative method!Jacobi} is characterized

 by a slow convergence\index{convergence} rate compared to some iterative methods\index{iterative method}
 (for example, Gauss-Seidel method\index{iterative method!Gauss-Seidel}). So, in this
 section, we present some solutions to reduce the execution time and the number of
 by a slow convergence\index{convergence} rate compared to some iterative methods\index{iterative method}
 (for example, Gauss-Seidel method\index{iterative method!Gauss-Seidel}). So, in this
 section, we present some solutions to reduce the execution time and the number of


@@ -776,7 +778,7 @@ vector elements leads to using twice the initial number of memory transactions.
 we apply the point red-black ordering\index{iterative method!red-black ordering}
 accordingly to the $y$-coordinate, as is shown in Figure~\ref{ch13:fig:06.02}. In
 this case, the vector elements having even $y$-coordinate are computed in parallel
 we apply the point red-black ordering\index{iterative method!red-black ordering}
 accordingly to the $y$-coordinate, as is shown in Figure~\ref{ch13:fig:06.02}. In
 this case, the vector elements having even $y$-coordinate are computed in parallel

-using the values of those having odd $y$-coordinate and then viceversa. Moreover,
+using the values of those having odd $y$-coordinate and then vice-versa. Moreover,

 in the GPU implementation of the parallel projected Richardson method (Section~\ref{ch13:sec:04}),
 we have shown that a subproblem of size $(NX\times ny\times nz)$ is decomposed into
 $nz$ grids of size $(NX\times ny)$. Then, each kernel is executed in parallel by
 in the GPU implementation of the parallel projected Richardson method (Section~\ref{ch13:sec:04}),
 we have shown that a subproblem of size $(NX\times ny\times nz)$ is decomposed into
 $nz$ grids of size $(NX\times ny)$. Then, each kernel is executed in parallel by


@@ -860,9 +862,9 @@ number of computing nodes on the cluster. This is due to the fact that the ratio
 between the time of the computation over that of the communication is reduced when
 the computations are performed on GPUs. Indeed, GPUs compute faster than CPUs and
 communications are more time-consuming. In this context, asynchronous algorithms
 between the time of the computation over that of the communication is reduced when
 the computations are performed on GPUs. Indeed, GPUs compute faster than CPUs and
 communications are more time-consuming. In this context, asynchronous algorithms

-are more scalable than synchronous ones. So, with large scale GPU clusters, synchronous\index{synchronous}
+are more scalable than synchronous ones. So, with large scale GPU clusters, synchronous\index{synchronous iterations}

 algorithms might be more penalized by communications, as can be deduced from Figure~\ref{ch13:fig:07}.
 algorithms might be more penalized by communications, as can be deduced from Figure~\ref{ch13:fig:07}.

-That is why we think that asynchronous\index{asynchronous} iterative algorithms
+That is why we think that asynchronous\index{asynchronous iterations} iterative algorithms

 are all the more interesting in this case.
 
 
 are all the more interesting in this case.