computation/communication overlapping}
\label{ch6:part2}
-In the previous part, we have seen how to efficiently implement overlap of
-computations (CPU and GPU) with communications (GPU transfers and inter-node
+In the previous section, we have seen how to efficiently implement overlap of
+computations (CPU and GPU) with communications (GPU transfers and internode
communications). However, we have previously shown that for some parallel
iterative algorithms, it is sometimes even more efficient to use an asynchronous
-scheme of iterations\index{iterations!asynchronous} \cite{HPCS2002,ParCo05,Para10}. In that case, the nodes do
-not wait for each others but they perform their iterations using the last
+scheme of iterations\index{asynchronous iterations} \cite{HPCS2002,ParCo05,Para10}. In that case, the nodes do
+not wait for each other but they perform their iterations using the last
external data they have received from the other nodes, even if this
data was produced \emph{before} the previous iteration on the other nodes.
Formally, if we denote by $f=(f_1,...,f_n)$ the function representing the
iterative process and by $x^t=(x_1^t,...,x_n^t)$ the values of the $n$ elements of
the system at iteration $t$, we pass from a synchronous iterative scheme of the
-form:
+form given in \Alg{algo:ch6p2sync}
%% \begin{algorithm}[H]
%% \caption{Synchronous iterative scheme}\label{algo:ch6p2sync}
%% \begin{Algo}
%% \end{Algo}
%% \end{algorithm}
\begin{algorithm}[H]
- \caption{Synchronous iterative scheme}\label{algo:ch6p2sync}
+ \caption{synchronous iterative scheme}\label{algo:ch6p2sync}
$x^{0}=(x_{1}^{0},...,x_{n}^{0})$\;
\For{ $t=0,1,...$} {
\For{ $i=1,...,n$}{
}
\end{algorithm}
-
\noindent
-to an asynchronous iterative scheme of the form:\\
+to an asynchronous iterative scheme of the form given in \Alg{algo:ch6p2async}.\\
%% \begin{algorithm}[H]
%% \caption{Asynchronous iterative scheme}\label{algo:ch6p2async}
%% \begin{Algo}
%% \end{Algo}
%% \end{algorithm}
\begin{algorithm}[H]
- \caption{Asynchronous iterative scheme}\label{algo:ch6p2async}
+ \caption{asynchronous iterative scheme}\label{algo:ch6p2async}
$x^{0}=(x_{1}^{0},...,x_{n}^{0})$\;
\For {$t=0,1,...$} {
\For{ $i=1,...,n$} {
}
}
\end{algorithm}
-where $s_j^i(t)$ is the iteration number of the production of the value $x_j$ of
-element $j$ that is used on element $i$ at iteration $t$ (see for example~\cite{BT89,
+In this scheme, $s_j^i(t)$ is the iteration number of the production of the value $x_j$ of
+element $j$ that is used on element $i$ at iteration $t$ (see, for example,~\cite{BT89,
FS2000} for further details).
Such schemes are called AIAC\index{AIAC} for \emph{Asynchronous Iterations and
Asynchronous Communications}. They combine two aspects that are respectively
if a node receives newer data only every four or five local iterations, it is
strongly probable that the evolution of its local iterative process will be
slower than if it receives data at every iteration. The key point here is that
-this frequency does not only depend on the hardware configuration of the
+not only does this frequency depend on the hardware configuration of the
parallel system but it also depends on the software that is used to
implement the algorithmic scheme.
algorithms has already been investigated in~\cite{SuperCo05}. Although the
features required to efficiently implement asynchronous schemes have not
changed, the available programming environments and computing hardware have
-evolved, in particular now GPUs are available. So, there is a need to reconsider the
+evolved, in particular now that GPUs are available. So, there is a need to reconsider the
implementation schemes of AIAC according to the new de facto standards for parallel
programming (communications and threads) as well as the integration of the GPUs.
One of the main objective here is to obtain a maximal overlap between the
-activities of the three types of devices that are the CPU, the GPU and the
+activities of the three types of devices: the CPU, the GPU, and the
network. Moreover, another objective is to present what we think
is the best compromise between the simplicity of the implementation and its
maintainability on one side and its
mechanism on top of our asynchronous scheme that can be used either statically or
dynamically during the application execution.
-Although there exist several programming environments for inter-node
-communications, multi-threading and GPU programming, a few of them have
-become \emph{de facto standards}, either due to their good stability, their ease
-of use and/or their wide adoption by the scientific community.
-Therefore, as in the previous section all the schemes presented in the following use MPI~\cite{MPI},
-OpenMP~\cite{openMP} and CUDA~\cite{CUDA}. However, there is no loss of
-generality as those schemes may easily be implemented with other libraries.
+Although there exist several programming environments for internode
+communications, multithreading, and GPU programming, a few of them have
+become \emph{de facto standards}, due to their good stability, their ease
+of use, and/or their wide adoption by the scientific community.
+Therefore, as in the previous section, all the schemes presented in the following use MPI~\cite{MPI},
+OpenMP~\cite{openMP}, and CUDA~\cite{CUDA}. However, there is no loss of
+generality as these schemes may easily be implemented with other libraries.
Finally, in order to stay as clear as possible, only the parts of code and
-variables related to the control of parallelism (communications, threads,...)
+variables related to the control of parallelism (communications, threads, etc.)
are presented in our schemes. The inner organization of data is not detailed as
it depends on the application. We only consider that we have two data arrays
(previous version and current version) and communication buffers. However, in
The first step toward our complete scheme is to implement a basic asynchronous
scheme that includes an actual overlap of the communications with the
computations\index{overlap!computation and communication}. In order to ensure that the communications are actually performed
-in parallel of the computations, it is necessary to use different threads. It
+in parallel with the computations, it is necessary to use different threads. It
is important to remember that asynchronous communications provided in
-communication libraries like MPI are not systematically performed in parallel of
+communication libraries such as MPI are not systematically performed in parallel with
the computations~\cite{ChVCV13,Hoefler08a}. So, the logical and classical way
to implement such an overlap is to use three threads: one for
-computing, one for sending and one for receiving. Moreover, since
-the communication is performed by threads, blocking synchronous communications\index{MPI!communication!blocking}\index{MPI!communication!synchronous}
+computing, one for sending, and one for receiving. Moreover, since
+the communication is performed by threads, blocking synchronous communications\index{MPI!blocking}\index{MPI!synchronous}
can be used without deteriorating the overall performance.
In this basic version, the termination\index{termination} of the global process is performed
-individually on each node according to their own termination. This can be guided by either a
+individually on each node according to its own termination. This can be guided by either a
number of iterations or a local convergence detection\index{convergence detection}. The important step at
the end of the process is to perform the receptions of all pending
communications in order to ensure the termination of the two communication
% \begin{algorithm}[H]
% \caption{Initialization of the basic asynchronous scheme.}
% \label{algo:ch6p2BasicAsync}
-\begin{Listing}{algo:ch6p2BasicAsync}{Initialization of the basic asynchronous scheme}
+\begin{Listing}{algo:ch6p2BasicAsync}{initialization of the basic asynchronous scheme}
// Variables declaration and initialization
// Controls the sendings from the computing thread
omp_lock_t lockSend;
In this scheme, the \texttt{lockRec} mutex\index{OpenMP!mutex} is not mandatory.
It is only used to ensure that data dependencies are actually exchanged at the
first iteration of the process. Data initialization and distribution
-(lines~16-17) are not detailed here because they are directly related to the
+(lines~20--21) are not detailed here because they are directly related to the
application. The important point is that, in most cases, they should be done
before the iterative process. The computing function is given in
\Lst{algo:ch6p2BasicAsyncComp}.
%\begin{algorithm}[H]
% \caption{Computing function in the basic asynchronous scheme.}
% \label{algo:ch6p2BasicAsyncComp}
-\begin{Listing}{algo:ch6p2BasicAsyncComp}{Computing function in the basic asynchronous scheme}
+\begin{Listing}{algo:ch6p2BasicAsyncComp}{computing function in the basic asynchronous scheme}
// Variables declaration and initialization
int iter = 1; // Number of the current iteration
double difference; // Variation of one element between two iterations
// Computation loop
while(!Finished){
- // Sendings of data dependencies if there is no previous sending
+ // Sending of data dependencies if there is no previous sending
// in progress
if(!SendsInProgress){
- // Potential copy of data to be sent in additional buffers
+ // Potential copy of data to be sent into additional buffers
...
// Change of sending state
SendsInProgress = 1;
\end{Listing}
%\end{algorithm}
-As mentioned above, it can be seen in line~18 of \Lst{algo:ch6p2BasicAsyncComp}
+As mentioned above, it can be seen in lines~19--21 of \Lst{algo:ch6p2BasicAsyncComp}
that the \texttt{lockRec} mutex is used only at the first iteration to wait for
the initial data dependencies before the computations. The residual\index{residual}, initialized
-in line~23 and computed in lines~34-37, is defined by the maximal difference
+in line~24 and computed in lines~35--38, is defined by the maximal difference
between the elements from two consecutive iterations. It is classically used to
detect the local convergence of the process on each node. In the more
complete schemes presented in the sequel, a global termination detection that
takes the states of all the nodes into account will be exhibited.
-Finally, the local convergence is tested and updated when necessary. In line~44,
+Finally, the local convergence is tested and updated when necessary. In line~45,
the \texttt{lockSend} mutex is unlocked to allow the sending function to send
final messages to the dependency nodes. Those messages are required to keep the
reception function alive until all the final messages have been received.
-Otherwise, a node could stop its reception function whereas other nodes are
+Otherwise, a node could stop its reception function while other nodes are
still trying to communicate with it. Moreover, a local sending of a final
-message to the node itself is required (line~45) to ensure that the reception
+message to the node itself is required (line~46) to ensure that the reception
function will not stay blocked in a message probing
-(see~\Lst{algo:ch6p2BasicAsyncReceptions}, line~11). This may happen if the node
-receives the final messages from its dependencies \emph{before} being itself in
+(see~\Lst{algo:ch6p2BasicAsyncReceptions}, line~12). This may happen if the node
+receives the final messages from its dependencies \emph{before} reaching its own
local convergence.
All the messages but this final local one are performed in the sending function
described in \Lst{algo:ch6p2BasicAsyncSendings}.
-
The main loop is only conditioned by the end of the computing process (line~4).
At each iteration, the thread waits for the permission from the computing thread
(according to the \texttt{lockSend} mutex). Then, data are sent with
%\begin{algorithm}[H]
% \caption{Sending function in the basic asynchronous scheme.}
% \label{algo:ch6p2BasicAsyncSendings}
-\begin{Listing}{algo:ch6p2BasicAsyncSendings}{Sending function in the basic asynchronous scheme}
+\begin{Listing}{algo:ch6p2BasicAsyncSendings}{sending function in the basic asynchronous scheme}
// Variables declaration and initialization
...
% \label{algo:ch6p2BasicAsyncReceptions}
\begin{Listing}{algo:ch6p2BasicAsyncReceptions}{Reception function in the basic asynchronous scheme}
// Variables declaration and initialization
-char countReceipts = 0; // Boolean indicating whether receptions are
-// counted or not
+char countReceipts = 1; // Boolean indicating whether receptions are
+ // counted or not
int nbEndMsg = 0; // Number of end messages received
-int arrived = 0; // Boolean indicating if a message is arrived
+int arrived = 0; // Boolean indicating if a message has arrived
int srcNd; // Source node of the message
int size; // Message size
}
// Reception of pending messages and counting of end messages
-do{ // Loop over the remaining incoming/waited messages
+do{ // Loop over the remaining incoming/end messages
MPI_Probe(MPI_ANY_SOURCE, MPI_ANY_TAG, MPI_COMM_WORLD, &status);
MPI_Get_count(&status, MPI_CHAR, &size);
// Actual reception in dummy buffer
\end{Listing}
%\end{algorithm}
-As in the sending function, the main loop of receptions is done while the
-iterative process is not \texttt{Finished}. In line~11, the thread waits until
-a message arrives on the node. Then, it performs the actual reception and the
-corresponding subsequent actions (potential data copies for data messages and
-counting for end messages). Lines 20-23 check that all data dependencies have
-been received before unlocking the \texttt{lockRec} mutex. As mentioned before,
-they are not mandatory and are included only to ensure that all data
-dependencies are received at the first iteration. Lines 25-28 are required to
-manage end messages that arrive on the node \emph{before} it reaches its own
-termination process. As the nodes are \emph{not} synchronized, this may happen.
-Finally, lines 34-43 perform the receptions of all pending communications,
-including the remaining end messages (at least the one from the node itself).
+As in the sending function, the main loop of receptions is done while the
+iterative process is not \texttt{Finished}. In line~12, the thread waits until
+a message arrives on the node. Then, it performs the actual reception and the
+corresponding subsequent actions (potential data copies for data messages and
+counting for end messages). Lines~23--26 check, only at the first iteration of
+computations, that all data dependencies have been received before unlocking the
+\texttt{lockRec} mutex. %As mentioned previously, they are not mandatory and are included only to
+Although this is not mandatory, it ensures that all data dependencies are
+received before starting the computations. % at the first iteration.
+Lines~28--31 are required to manage end messages that arrive on the
+node \emph{before} it reaches its own termination process. As the nodes are
+\emph{not} synchronized, this may happen. Finally, lines~37--46 perform the
+receptions of all pending communications, including the remaining end messages
+(at least the one from the node itself).
\medskip
-So, with those algorithms, we obtain a quite simple and efficient asynchronous
+So, with these algorithms, we obtain a quite simple and efficient asynchronous
iterative scheme. It is interesting to notice that GPU computing can be easily
included in the computing thread. This will be fully addressed in
-paragraph~\ref{ch6:p2GPUAsync}. However, before presenting the complete
+Section~\ref{ch6:p2GPUAsync}. However, before presenting the complete
asynchronous scheme with GPU computing, we have to detail how our initial scheme
can be made synchronous.
% code, which tends to improve the implementation and maintenance costs when both
% versions are required.
% The second one is that
-In our context, the interest of being able to dynamically change the operating
-mode (sync/async) during the process execution, is that this strongly simplifies
+In our context, being able to dynamically change the operating
+mode (sync/async) during the process execution strongly simplifies
the global convergence detection. In fact, our past experience in the design
and implementation of global convergence detection in asynchronous
-algorithms~\cite{SuperCo05, BCC07, Vecpar08a}, have led us to the conclusion
+algorithms~\cite{SuperCo05, BCC07, Vecpar08a} has led us to the conclusion
that although a decentralized detection scheme is possible and may be more
efficient in some situations, its much higher complexity is an
obstacle to actual use in practice, especially in industrial contexts where
during the execution of the process in order to check the global convergence.
This is why we need to synchronize our asynchronous scheme.
-In each algorithm of the initial scheme, we only exhibit the additional code
+In each algorithm of the initial scheme, we only give the additional code
required to change the operating mode.
%\begin{algorithm}[H]
% \caption{Initialization of the synchronized scheme.}
% \label{algo:ch6p2Sync}
-\begin{Listing}{algo:ch6p2Sync}{Initialization of the synchronized scheme}
+\begin{Listing}{algo:ch6p2Sync}{initialization of the synchronized scheme}
// Variables declarations and initialization
...
// Controls the synchronous exchange of local states
in \Lst{algo:ch6p2SyncComp}, those messages contain only a boolean indicating
for each node if it is in local convergence\index{convergence!local}. So, once all the states are
received on a node, it is possible to determine if all the nodes are in local
-convergence, and thus to detect the global convergence. The \texttt{lockIter}
+convergence and, thus, to detect the global convergence. The \texttt{lockIter}
mutex is used to synchronize all the nodes at the end of each iteration. There
-are also two new variables that respectively represent the local state of the
+are also two new variables that represent the local state of the
node (\texttt{localCV}) according to the iterative process (convergence) and the
number of other nodes that are in local convergence (\texttt{nbOtherCVs}).
%\begin{algorithm}[H]
% \caption{Computing function in the synchronized scheme.}
% \label{algo:ch6p2SyncComp}
-\begin{Listing}{algo:ch6p2SyncComp}{Computing function in the synchronized scheme}
+\begin{Listing}{algo:ch6p2SyncComp}{computing function in the synchronized scheme}
// Variables declarations and initialization
...
// Computation loop
while(!Finished){
- // Sendings of data dependencies at @\emph{each}@ iteration
+ // Sending of data dependencies at @\color{white}\emph{\textbf{each}}@ iteration
// Potential copy of data to be sent in additional buffers
...
omp_unset_lock(&lockSend);
- // Blocking receptions at @\emph{each}@ iteration
+ // Blocking receptions at @\color{white}\emph{\textbf{each}}@ iteration
omp_set_lock(&lockRec);
// Local computation
Between lines~6 and~7, the use of the flag \texttt{SendsInProgress} is no longer
needed since the sends are performed at each iteration. In line~12, the
thread waits for the data receptions from its dependencies. In
-lines~26-34, the local states are determined and exchanged among all nodes. A
+lines~27--34, the local states are determined and exchanged among all nodes. A
new message tag (\texttt{tagState}) is required for identifying those messages.
In line~37, the global termination state is determined. When it is reached,
-lines~38-42 change the \texttt{Finished} boolean to stop the iterative process,
+lines~39--42 change the \texttt{Finished} boolean to stop the iterative process
and send the end messages. Otherwise each node resets its local state
-information about the other nodes and a global barrier is done between all the
+information about the other nodes and a global barrier is added between all the
nodes at the end of each iteration with another new tag (\texttt{tagIter}).
That barrier is needed to ensure that data messages from successive iterations
are actually received during the \emph{same} iteration on the destination nodes.
Nevertheless, it is not useful at the termination of the global process as it is
replaced by the global exchange of end messages.
-There is no big modification induced by the synchronization in the sending
-function. The only change could be the suppression of line~11 that is not useful
-in this case. Apart from that, the function stays the same as
-in \Lst{algo:ch6p2BasicAsyncSendings}.
+There is no big modification induced by the synchronization in the sending
+function. The function stays almost the same as in
+\Lst{algo:ch6p2BasicAsyncSendings}. The only change could be the suppression of
+line~11 that is not useful in this case.
In the reception function, given in \Lst{algo:ch6p2SyncReceptions}, there are
-mainly two insertions (in lines~19-30 and 31-40), corresponding to the
+mainly two insertions (in lines~19--31 and 32--42), corresponding to the
additional types of messages to receive. There is also the insertion of three
variables that are used for the receptions of the new message types. In
-lines~24-29 and 34-39 are located messages counting and mutex unlocking
+lines~24--30 and 35--41 are located messages counting and mutex unlocking
mechanisms that are used to block the computing thread at the corresponding steps of its
execution. They are similar to the mechanism used for managing the end messages
at the end of the entire process. Line~23 directly updates the
number of other nodes that are in local convergence by adding the
received state of the source node. This is possible due to the encoding that is used to
-represent the local convergence (1) and the non-convergence (0).
+represent the local convergence (1) and the nonconvergence (0).
%\begin{algorithm}[H]
% \caption{Reception function in the synchronized scheme.}
% \label{algo:ch6p2SyncReceptions}
-\begin{Listing}{algo:ch6p2SyncReceptions}{Reception function in the synchronized scheme}
+\begin{Listing}{algo:ch6p2SyncReceptions}{reception function in the synchronized scheme}
// Variables declarations and initialization
...
int nbStateMsg = 0; // Number of local state messages received
case tagState: // Management of local state messages
// Actual reception of the message
MPI_Recv(&recvdState, 1, MPI_CHAR, status.MPI_SOURCE, tagState, MPI_COMM_WORLD, &status);
- // Updates of numbers of stabilized nodes and received state msgs
+ // Updates of numbers of stabilized nodes and recvd state msgs
nbOtherCVs += recvdState;
nbStateMsg++;
// Unlocking of the computing thread when states of all other
case tagIter: // Management of "end of iteration" messages
// Actual reception of the message in dummy buffer
MPI_Recv(dummyBuffer, 1, MPI_CHAR, status.MPI_SOURCE, tagIter, MPI_COMM_WORLD, &status);
- nbIterMsg++; // Update of the nb of iteration messages
+ nbIterMsg++; // Update of the number of iteration messages
// Unlocking of the computing thread when iteration messages
// are received from all other nodes
if(nbIterMsg == nbP - 1){
}
// Reception of pending messages and counting of end messages
-do{ // Loop over the remaining incoming/waited messages
+do{ // Loop over the remaining incoming/end messages
...
}while(arrived == 1 || nbEndMsg < nbDeps + 1);
\end{Listing}
Now that we can synchronize our asynchronous scheme, the final step is to
dynamically alternate the two operating modes in order to regularly check the
global convergence of the iterative process. This is detailed in the following
-paragraph together with the inclusion of GPU computing in the final asynchronous
+section together with the inclusion of GPU computing in the final asynchronous
scheme.
-\subsection{Asynchronous scheme using MPI, OpenMP and CUDA}
+\subsection{Asynchronous scheme using MPI, OpenMP, and CUDA}
\label{ch6:p2GPUAsync}
As mentioned above, the strategy proposed to obtain a good compromise between
The last problem is to decide \emph{when} to switch from the
asynchronous to the synchronous mode. Here again, for the sake of simplicity, any
-asynchronous mechanism for \emph{detecting} such instant is avoided and we
+asynchronous mechanism for \emph{detecting} such moment is avoided, and we
prefer to use a mechanism that is local to each node. Obviously, that local system must
rely neither on the number of local iterations done nor on the local
convergence. The former would slow down the fastest nodes according to the
slowest ones. The latter would provoke too much synchronization because the
-residuals on all nodes commonly do not evolve in the same way and, in most
+residuals on all nodes generally do not evolve in the same way, and in most
cases, there is a convergence wave phenomenon throughout the elements. So, a
good solution is to insert a local timer mechanism on each node with a given
initial duration. Then, that duration may be modified during the execution
according to the successive results of the synchronous sections.
Another problem induced by entering synchronous mode from the asynchronous one
-is the possibility to receive some data messages
+is the possibility of receiving some data messages
from previous asynchronous iterations during synchronous iterations. This could lead to deadlocks. In order
-to avoid this, a wait of the end of previous send is added to the
+to avoid this, a wait for the end of previous send is added to the
transition between the two modes. This is implemented by replacing the variable
-\texttt{SendsInProgress} by a mutex \texttt{lockSendsDone} which is unlocked
+\texttt{SendsInProgress} with a mutex \texttt{lockSendsDone} which is unlocked
once all the messages have been sent in the sending function. Moreover, it is
also necessary to stamp data messages\index{message!stamping} (by the function \texttt{stampData}) with
-a Boolean indicating whether they have been sent during a synchronous or
+a boolean indicating whether they have been sent during a synchronous or
asynchronous iteration. Then, the \texttt{lockRec} mutex is unlocked only
after to the complete reception of data messages from synchronous
iterations. The message ordering of point-to-point communications in MPI
-together with the barrier at the end of each iteration ensure two important
+and the barrier at the end of each iteration ensure two important
properties of this mechanism. First, data messages from previous
asynchronous iterations will be received but not taken into account during
-synchronous sections. Then, a data message from a synchronous
-iteration cannot be received in another synchronous iteration. In the
+synchronous sections. Then, a data message from a given synchronous
+iteration cannot be received during another synchronous iteration. In the
asynchronous sections, no additional mechanism is needed as there are no such
constraints concerning the data receptions.
are mainly related to the computing thread. Small additions or modifications
are also required in the main process and the other threads.
-In the main process, two new variables are added to store respectively the main
+In the main process, two new variables are added to store the main
operating mode of the iterative process (\texttt{mainMode}) and the duration of
asynchronous sections (\texttt{asyncDuration}). Those variables are
-initialized by the programmer. The \texttt{lockSendsDone} mutex is also declared,
-initialized (locked) and destroyed with the other mutex in this process.
+initialized by the programmer. The mutex \texttt{lockSendsDone} is also declared,
+initialized (locked), and destroyed with the other mutex in this process.
In the computing function, shown in \Lst{algo:ch6p2AsyncSyncComp}, the
-modifications consist of the insertion of the timer mechanism and of the conditions
+modifications consist of the insertion of the timer mechanism and the tests
to differentiate the actions to be done in each mode. Some additional variables
are also required to store the current operating mode in action during the
execution (\texttt{curMode}), the starting time of the current asynchronous
-section (\texttt{asyncStart}) and the number of successive synchronous
+section (\texttt{asyncStart}), and the number of successive synchronous
iterations done (\texttt{nbSyncIter}).
%\begin{algorithm}[H]
% \caption{Computing function in the final asynchronous scheme.}% without GPU computing.}
% \label{algo:ch6p2AsyncSyncComp}
%\pagebreak
-\begin{Listing}{algo:ch6p2AsyncSyncComp}{Computing function in the final asynchronous scheme}% without GPU computing.}
+\begin{Listing}{algo:ch6p2AsyncSyncComp}{computing function in the final asynchronous scheme}% without GPU computing.}
// Variables declarations and initialization
...
OpMode curMode = SYNC;// Current operating mode (always begin in sync)
}
}
- // Sendings of data dependencies
+ // Sending of data dependencies
if(curMode == SYNC || !SendsInProgress){
...
}
}
// Local computation
- // (init of residual, arrays swapping and iteration computation)
+ // (init of residual, arrays swapping, and iteration computation)
...
- // Checking of convergences (local & global) only in sync mode
+ // Checking convergences (local & global) only in sync mode
if(curMode == SYNC){
// Local convergence checking (residual under threshold)
...
// Blocking global exchange of local states of the nodes
...
// Determination of global convergence (all nodes in local CV)
- // Stop of the iterative process and sending of end messages
- // or Re-initialization of state information and iteration barrier
+ // Stopping the iterative process and sending end messages
+ // or reinitialization of state information and iteration barrier
...
}
}
%\end{algorithm}
In the sending function, the only modification is the replacement in line~11 of
-the assignment of variable \texttt{SendsInProgress} by the unlocking of
+the assignment of variable \texttt{SendsInProgress} with the unlocking of
\texttt{lockSendsDone}. Finally, in the reception function, the only
-modification is the insertion before line~19
+modification is the insertion before line~21
of \Lst{algo:ch6p2BasicAsyncReceptions} of the extraction of the stamp from the
message and its counting among the receipts only if the stamp is \texttt{SYNC}.
The final step to get our complete scheme using GPU is to insert the GPU
management in the computing thread. The first possibility, detailed
in \Lst{algo:ch6p2syncGPU}, is to simply replace the
-CPU kernel (lines~41-43 in \Lst{algo:ch6p2AsyncSyncComp}) by a blocking GPU kernel call. This includes data
+CPU kernel (lines~42--44 in \Lst{algo:ch6p2AsyncSyncComp}) by a blocking GPU kernel call. This includes data
transfers from the node RAM to the GPU RAM, the launching of the GPU kernel, the
-waiting for kernel completion and the results transfers from GPU RAM to
+waiting for kernel completion, and the results transfers from GPU RAM to
node RAM.
%\begin{algorithm}[H]
% \caption{Computing function in the final asynchronous scheme.}
% \label{algo:ch6p2syncGPU}
-\begin{Listing}{algo:ch6p2syncGPU}{Computing function in the final asynchronous scheme}
+\begin{Listing}{algo:ch6p2syncGPU}{computing function in the final asynchronous scheme}
// Variables declarations and initialization
...
dim3 Dg, Db; // CUDA kernel grids
// Computation loop
while(!Finished){
// Determination of the dynamic operating mode, sendings of data
- // dependencies and blocking data receptions in sync mode
+ // dependencies, and blocking data receptions in sync mode
...
// Local GPU computation
// Data transfers from node RAM to GPU
CHECK_CUDA_SUCCESS(cudaMemcpyToSymbol(dataOnGPU, dataInRAM, inputsSize, 0, cudaMemcpyHostToDevice), "Data transfer");
... // There may be several data transfers: typically A and b in
- // linear problems
+ // linear problems of the form A.x = b
// GPU grid definition
Db.x = BLOCK_SIZE_X; // BLOCK_SIZE_# are kernel design dependent
Db.y = BLOCK_SIZE_Y;
%\end{algorithm}
This scheme provides asynchronism through a cluster of GPUs as well as a
-complete overlap of communications with GPU computations (similarly
-to~\Sec{ch6:part1}). However, the autonomy of GPU devices according to their
+complete overlap of communications with GPU computations (similar
+to the one described in~\Sec{ch6:part1}). However, the autonomy of GPU devices according to their
host can be further exploited in order to perform some computations on the CPU
while the GPU kernel is running. The nature of computations that can be done by
the CPU may vary depending on the application. For example, when processing data
-streams (pipelines), pre-processing of next data item and/or post-processing of
-previous result can be done on the CPU while the GPU is processing the current
+streams (pipelines), pre-processing of the next data item and/or post-processing
+of the previous result can be done on the CPU while the GPU is processing the current
data item. In other cases, the CPU can perform \emph{auxiliary}
-computations\index{computation!auxiliary}
+computations\index{computation auxiliary}
that are not absolutely required to obtain the result but that may accelerate
the entire iterative process. Another possibility would be to distribute the
main computations between the GPU and CPU. However, this
-usually leads to poor performance increases. This is mainly due to data
+usually leads to poor performance increases mainly due to data
dependencies that often require additional transfers between CPU and GPU.
So, if we consider that the application enables such overlap of
computations, its implementation is straightforward as it consists in inserting
-the additional CPU computations between lines~23 and~24
-in \Lst{algo:ch6p2syncGPU}. Nevertheless, such scheme is fully efficient only
+the additional CPU computations between lines~25 and~26
+in \Lst{algo:ch6p2syncGPU}. Nevertheless, such a scheme is fully efficient only
if the computation times on both sides are similar.
In some cases, especially with auxiliary computations, another interesting
solution is to add a fourth CPU thread to perform them. This suppresses the
duration constraint over those optional computations as they are performed in
-parallel of the main iterative process, without blocking it. Moreover, this
+parallel with the main iterative process, without blocking it. Moreover, this
scheme stays coherent with current architectures as most nodes include four CPU
cores. The algorithmic scheme of such context of complete overlap of
CPU/GPU computations and communications is described in
-Listings~\ref{algo:ch6p2FullOverAsyncMain},~\ref{algo:ch6p2FullOverAsyncComp1}
-and~\ref{algo:ch6p2FullOverAsyncComp2}, where we suppose that auxiliary
+Listings~\ref{algo:ch6p2FullOverAsyncMain},~\ref{algo:ch6p2FullOverAsyncComp1},
+and~\ref{algo:ch6p2FullOverAsyncComp2}, where we assume that auxiliary
computations use intermediate results of the main computation process from any previous iteration. This may be
different according to the application.
%\begin{algorithm}[H]
% \caption{Initialization of the main process of complete overlap with asynchronism.}
% \label{algo:ch6p2FullOverAsyncMain}
-\pagebreak
-\begin{Listing}{algo:ch6p2FullOverAsyncMain}{Initialization of the main process of complete overlap with asynchronism}
+%\pagebreak
+\begin{Listing}{algo:ch6p2FullOverAsyncMain}{initialization of the main process of complete overlap with asynchronism}
// Variables declarations and initialization
...
omp_lock_t lockAux; // Informs main thread about new aux results
omp_lock_t lockWrite; // Controls exclusion of results access
... auxRes ... ; // Results of auxiliary computations
-// Parameters reading, MPI initialization, data initialization and
+// Parameters reading, MPI initialization, and data initialization and
// distribution
...
// OpenMP initialization
%\begin{algorithm}[H]
% \caption{Computing function in the final asynchronous scheme with CPU/GPU overlap.}
% \label{algo:ch6p2FullOverAsyncComp1}
-\pagebreak
-\begin{Listing}{algo:ch6p2FullOverAsyncComp1}{Computing function in the final asynchronous scheme with CPU/GPU overlap}
+%\pagebreak
+\begin{Listing}{algo:ch6p2FullOverAsyncComp1}{computing function in the final asynchronous scheme with CPU/GPU overlap}
// Variables declarations and initialization
...
dim3 Dg, Db; // CUDA kernel grids
// Computation loop
while(!Finished){
- // Determination of the dynamic operating mode, sendings of data
- // dependencies and blocking data receptions in sync mode
+ // Determination of the dynamic operating mode, sending of data
+ // dependencies, and blocking data receptions in sync mode
...
// Local GPU computation
- // Data transfers from node RAM to GPU, GPU grid definition and init
- // of shared mem
+ // Data transfers from node RAM to GPU, GPU grid definition,
+ // and init of shared memory
CHECK_CUDA_SUCCESS(cudaMemcpyToSymbol(dataOnGPU, dataInRAM, inputsSize, 0, cudaMemcpyHostToDevice), "Data transfer");
...
// Kernel call
gpuKernelName<<<Dg,Db>>>(... @\emph{kernel parameters}@ ...);
- // Potential pre/post-treatments in pipeline like computations
+ // Potential pre-/post-treatments in pipeline-like computations
...
// Waiting for kernel completion
cudaDeviceSynchronize();
// Results transfer from GPU to node RAM
omp_set_lock(&lockWrite); // Wait for write access to resultsInRam
CHECK_CUDA_SUCCESS(cudaMemcpyFromSymbol(resultsInRam, resultsOnGPU, resultsSize, 0, cudaMemcpyDeviceToHost), "Results transfer");
- // Potential post-treatments in non-pipeline computations
+ // Potential post-treatments in non pipeline computations
...
omp_unset_lock(&lockWrite); // Give back read access to aux thread
omp_test_lock(&lockRes);
...
// Determination of global convergence (all nodes in local CV)
if(cvLocale == 1 && nbCVLocales == nbP-1){
- // Stop of the iterative process and sending of end messages
+ // Stopping the iterative process and sending end messages
...
- // Unlocking of aux thread for termination
+ // Unlocking aux thread for termination
omp_test_lock(&lockRes);
omp_unset_lock(&lockRes);
}else{
- // Re-initialization of state information and iteration barrier
+ // Reinitialization of state information and iteration barrier
...
}
}
%\begin{algorithm}[H]
% \caption{Auxiliary computing function in the final asynchronous scheme with CPU/GPU overlap.}
% \label{algo:ch6p2FullOverAsyncComp2}
-\pagebreak
-\begin{Listing}{algo:ch6p2FullOverAsyncComp2}{Auxiliary computing function in the final asynchronous scheme with CPU/GPU overlap}
+%\pagebreak
+\begin{Listing}{algo:ch6p2FullOverAsyncComp2}{auxiliary computing function in the final asynchronous scheme with CPU/GPU overlap}
// Variables declarations and initialization
... auxInput ... // Local array for input data
%\end{algorithm}
As can be seen in \Lst{algo:ch6p2FullOverAsyncMain}, there are three additional
-mutex (\texttt{lockAux}, \texttt{lockRes} and \texttt{lockWrite}) that are used
-respectively to inform the main computation thread that new auxiliary results
-are available (lines~20-21 in \Lst{algo:ch6p2FullOverAsyncComp2} and line~29 in
+mutex (\texttt{lockAux}, \texttt{lockRes}, and \texttt{lockWrite}) that are used
+to inform the main computation thread that new auxiliary results
+are available (lines~20--21 in \Lst{algo:ch6p2FullOverAsyncComp2} and line~31 in
\Lst{algo:ch6p2FullOverAsyncComp1}), to inform the auxiliary thread that new
-results from the main thread are available (lines~25-26
+results from the main thread are available (lines~27--28
in \Lst{algo:ch6p2FullOverAsyncComp1} and line~7
in \Lst{algo:ch6p2FullOverAsyncComp2}), and to perform exclusive accesses to the
-results from those two threads (lines~20,~24
-in \Lst{algo:ch6p2FullOverAsyncComp1} and 9,~13
+results from those two threads (lines~22 and 26
+in \Lst{algo:ch6p2FullOverAsyncComp1} and 9 and 13
in \Lst{algo:ch6p2FullOverAsyncComp2}). Also, an additional array
(\texttt{auxRes}) is required to store the results of the auxiliary computations
as well as a local array for the input of the auxiliary function
(\texttt{auxInput}). That last function has the same general organization as the
-send/receive ones, that is a global loop conditioned by the end of the global
+send/receive ones, that is, a global loop conditioned by the end of the global
process. At each iteration in this function, the thread waits for the
availability of new results produced by the main computation thread. This avoids
-to perform the same computations several times with the same input data.
+ performing the same computations several times with the same input data.
Then, input data of auxiliary computations
% (as supposed here, they often
% correspond to the results of the main computations, but may sometimes be
computations are performed. When they are completed, the associated mutex is
unlocked to signal the availability of those auxiliary results to the main
computing thread. The main thread regularly checks this availability at the end
-of its iterations and takes them into account whenever this is possible.
+of its iterations and takes them into account whenever possible.
-Finally, we obtain an algorithmic scheme allowing maximal overlap between
-CPU and GPU computations as well as communications. It is worth noticing that
-such scheme is also usable for systems without GPUs but 4-cores nodes.
+Finally, we obtain an algorithmic scheme allowing maximal overlap between CPU
+and GPU computations as well as communications. It is worth noticing that such
+scheme is also efficiently usable for systems without GPUs but with nodes having
+at least four cores. In such contexts, each thread in \Lst{algo:ch6p2FullOverAsyncMain} can
+be executed on distinct cores.
\subsection{Experimental validation}
\label{sec:ch6p2expes}
As in~\Sec{ch6:part1}, we validate the feasibility of our asynchronous scheme
with some experiments performed with a representative example of scientific
-application. It is a three-dimensional version of the
-advection-diffusion-reaction process\index{PDE example} that models the evolution of the
+application. This three-dimensional version of the
+advection-diffusion-reaction process\index{PDE example} models the evolution of the
concentrations of two chemical species in shallow waters. As this process is
dynamic in time, the simulation is performed for a given number of consecutive
time steps. This implies two nested loops in the iterative process, the outer
one for the time steps and the inner one for solving the problem at each time.
-Full details about this PDE problem can be found in~\cite{ChapNRJ2011}. That
+Full details about this PDE problem can be found in~\cite{ChapNRJ2011}. This
two-stage iterative process implies a few adaptations of the general scheme
presented above in order to include the outer iterations over the time steps, but the
inner iterative process closely follows the same scheme.
We show two series of experiments performed with 16 nodes of the first cluster
described in~\Sec{ch6:p1expes}. The first one deals with the comparison of
synchronous and asynchronous computations. The second one is related to the use
-of auxiliary computations. In the context of our PDE application, they consist in
+of auxiliary computations. In the context of our PDE application, they consist of
the update of the Jacobian of the system.
\subsubsection*{Synchronous and asynchronous computations}
\label{fig:ch6p2syncasync}
\end{figure}
-The results obtained show that the asynchronous version is sensibly faster than
+The results obtained show that the asynchronous version is significantly faster than
the synchronous one for smaller problem sizes, then it becomes similar or even
a bit slower for larger problem sizes. A closer comparison of computation and
-communication times in each execution confirms that this behavior is consistent.
+communication times of each execution confirms that this behavior is consistent.
The asynchronous version is interesting if communication time
is similar or larger than computation time. In our example, this is the
case up to a problem size between 50 and 60. Then, computations become longer
$F'\times \Delta x=-F$ to update $x$ with $\Delta x$. There are two levels of
iterations, the inner level to get a stabilized version of $x$, and the outer
level to compute $x$ at the successive time steps in the simulation process. In
-this context, classical algorithms either compute $F'$ only at the first iteration
+this context, classic algorithms either compute $F'$ at only the first iteration
of each time step or at some iterations but not all because the computation of $F'$ is done
in the main iterative process and it has a relatively high computing cost.
However, with the scheme presented above, it is possible to continuously compute
-new versions of $F'$ in parallel to the main iterative process without
+new versions of $F'$ in parallel with the main iterative process without
penalizing it. Hence, $F'$ is updated as often as possible and taken into
account in the main computations when it is relevant. So, the Newton process
should be accelerated a little bit.
We compare the performance obtained with overlapped Jacobian updatings and
-non-overlapped ones for several problem sizes, see~\Fig{fig:ch6p2aux}.
+nonoverlapped ones for several problem sizes (see~\Fig{fig:ch6p2aux}).
\begin{figure}[h]
\centering
\includegraphics[width=.75\columnwidth]{Chapters/chapter6/curves/recouvs.pdf}
\end{figure}
The overlap is clearly efficient as the computation times with overlapping
-Jacobian updatings are much better than without overlap. Moreover, the ratio
-between the two versions tend to increase with the problem size, which is as
+Jacobian updatings are much better than the ones without overlap. Moreover, the ratio
+between the two versions tends to increase with the problem size, which is as
expected. Also, we have tested the application without auxiliary computations at
all, that is, the Jacobian is computed only once at the beginning of each time
step of the simulation. The results for this last version are quite similar to
the overlapped auxiliary computations, and even better for small problem sizes.
-The fact that no sensible gain can be seen on this range of problem sizes is due
+The fact that no significant gain can be seen on this range of problem sizes is due
to the limited number of Jacobian updates taken into account in the main
computation. This happens when the Jacobian update is as long as
several iterations of the main process. So, the benefit is reduced in this
particular case.
-Those results show two things; first, auxiliary computations do not induce great
-overhead in the whole process. Second, for this particular application the
+Those results show two things. First, auxiliary computations do not induce great
+overhead in the whole process. Second, for this particular application the
choice of updating the Jacobian matrix as auxiliary computations does not speed
up the iterative process. This does not question the parallel scheme in itself
-but merely points out the difficulty to identify relevant auxiliary
+but merely points out the difficulty of identifying relevant auxiliary
computations. Indeed, this identification depends on the considered application
and requires a profound specialized analysis.
Another interesting choice could be the computation of load estimation for
dynamic load balancing, especially in decentralized diffusion strategies where
-loads are transferred between neighboring nodes~\cite{BCVG11}. In such case,
+loads are transferred between neighboring nodes~\cite{BCVG11}. In such a case,
the load evaluation and the comparison with other nodes can be done in parallel
-of the main computations without perturbing them.
+with the main computations without perturbing them.
%%% Local Variables:
%%% mode: latex
