X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/book_gpu.git/blobdiff_plain/637049bfdd22c413d65ad0548d2d18a70a1fa6be..55ce7168c6e69a2462d76c95dc9a5298ceedb04f:/BookGPU/Chapters/chapter6/PartieAsync.tex diff --git a/BookGPU/Chapters/chapter6/PartieAsync.tex b/BookGPU/Chapters/chapter6/PartieAsync.tex index 4a2b6d2..1617ffb 100644 --- a/BookGPU/Chapters/chapter6/PartieAsync.tex +++ b/BookGPU/Chapters/chapter6/PartieAsync.tex @@ -2,19 +2,19 @@ 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} @@ -27,7 +27,7 @@ form: %% \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$}{ @@ -36,9 +36,8 @@ form: } \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} @@ -56,7 +55,7 @@ to an asynchronous iterative scheme of the form:\\ %% \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$} { @@ -69,8 +68,8 @@ to an asynchronous iterative scheme of the form:\\ } } \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 @@ -87,7 +86,7 @@ which new data arrives on each node. Typically, 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. @@ -95,11 +94,11 @@ The impact of the programming environments used to implement asynchronous 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 @@ -112,16 +111,16 @@ with full overlap. Between these two extremes, we propose a synchronization 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 @@ -134,17 +133,17 @@ avoid data copies. 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 @@ -155,13 +154,17 @@ So, the global organization of this scheme is set up in \Lst{algo:ch6p2BasicAsyn % \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 -omp_lock_t lockSend; // Controls the sendings from the computing thread -omp_lock_t lockRec; // Ensures the initial reception of external data -char Finished = 0; // Boolean indicating the end of the process -char SendsInProgress = 0; // Boolean indicating if previous data sendings are still in progress -double Threshold; // Threshold of the residual for convergence detection +// Controls the sendings from the computing thread +omp_lock_t lockSend; +// Ensures the initial reception of external data +omp_lock_t lockRec; +char Finished = 0; // Boolean indicating the end of the process +// Boolean indicating if previous data sendings are still in progress +char SendsInProgress = 0; +// Threshold of the residual for convergence detection +double Threshold; // Parameters reading ... @@ -177,9 +180,10 @@ MPI_Comm_rank(MPI_COMM_WORLD, &numP); // OpenMP initialization (mainly declarations and setting up of locks) omp_set_num_threads(3); omp_init_lock(&lockSend); -omp_set_lock(&lockSend); // Initially locked, unlocked to start sendings +omp_set_lock(&lockSend);//Initially locked, unlocked to start sendings omp_init_lock(&lockRec); -omp_set_lock(&lockRec); // Initially locked, unlocked when initial data are received +//Initially locked, unlocked when initial data are received +omp_set_lock(&lockRec); #pragma omp parallel { @@ -220,7 +224,7 @@ MPI_Finalize(); 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}. @@ -228,7 +232,7 @@ before the iterative process. The computing function is given in %\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 @@ -236,9 +240,10 @@ double residual; // Residual of the current iteration // Computation loop while(!Finished){ - // Sendings of data dependencies if there is no previous sending in progress + // 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; @@ -282,30 +287,29 @@ while(!Finished){ \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 @@ -318,7 +322,7 @@ the main loop, the final messages are sent to the dependencies of the node. %\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 ... @@ -345,9 +349,10 @@ The last function, detailed in \Lst{algo:ch6p2BasicAsyncReceptions}, does all th % \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 @@ -359,10 +364,12 @@ while(!Finished){ // Management of data messages switch(status.MPI_TAG){ case tagCom: // Management of data messages - srcNd = status.MPI_SOURCE; // Get the source node of the message + // Get the source node of the message + srcNd = status.MPI_SOURCE; // Actual data reception in the corresponding buffer MPI_Recv(dataBufferOf(srcNd), nbDataOf(srcNd), dataTypeOf(srcNd), srcNd, tagCom, MPI_COMM_WORLD, &status); - // Unlocking of the computing thread when data are received from all dependencies + // Unlocking of the computing thread when data are received + // from all dependencies if(countReceipts == 1 && ... @\emph{receptions from ALL dependencies}@ ...){ omp_unset_lock(&lockRec); countReceipts = 0; // No more counting after first iteration @@ -377,7 +384,7 @@ while(!Finished){ } // 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 @@ -390,24 +397,26 @@ do{ // Loop over the remaining incoming/waited messages \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. @@ -429,11 +438,11 @@ threshold. % 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 @@ -445,19 +454,23 @@ consists in dynamically changing the operating mode between asynchronous and syn 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 ... -omp_lock_t lockStates; // Controls the synchronous exchange of local states -omp_lock_t lockIter; // Controls the synchronization at the end of each iteration -char localCV = 0; // Boolean indicating whether the local stabilization is reached or not -int nbOtherCVs = 0; // Number of other nodes being in local stabilization +// Controls the synchronous exchange of local states +omp_lock_t lockStates; +// Controls the synchronization at the end of each iteration +omp_lock_t lockIter; +//Boolean indicating whether the local stabilization is reached or not +char localCV = 0; +// Number of other nodes being in local stabilization +int nbOtherCVs = 0; // Parameters reading ... @@ -468,9 +481,12 @@ int nbOtherCVs = 0; // Number of other nodes being in local stabilization // OpenMP initialization (mainly declarations and setting up of locks) ... omp_init_lock(&lockStates); -omp_set_lock(&lockStates); // Initially locked, unlocked when all state messages are received +// Initially locked, unlocked when all state messages are received +omp_set_lock(&lockStates); omp_init_lock(&lockIter); -omp_set_lock(&lockIter); // Initially locked, unlocked when all "end of iteration" messages are received +// Initially locked, unlocked when all "end of iteration" messages are +// received +omp_set_lock(&lockIter); // Threads launching #pragma omp parallel @@ -500,9 +516,9 @@ receptions of all state messages coming from the other nodes. As shown 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}). @@ -512,18 +528,18 @@ in \Lst{algo:ch6p2SyncComp}. %\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 @@ -548,7 +564,7 @@ while(!Finished){ // Waiting for the state messages receptions from the other nodes omp_set_lock(&lockStates); - // Determination of global convergence (if all nodes are in local CV) + //Determination of global convergence (if all nodes are in local CV) if(localCV + nbOtherCVs == nbP){ // Entering global CV state Finished = 1; @@ -578,39 +594,39 @@ Most of the added code is related to the waiting for specific communications. 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 @@ -632,10 +648,11 @@ while(!Finished){ 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 nodes are received + // Unlocking of the computing thread when states of all other + // nodes are received if(nbStateMsg == nbP-1){ nbStateMsg = 0; omp_unset_lock(&lockStates); @@ -644,8 +661,9 @@ while(!Finished){ 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 - // Unlocking of the computing thread when iteration messages are received from all other nodes + 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){ nbIterMsg = 0; omp_unset_lock(&lockIter); @@ -656,7 +674,7 @@ while(!Finished){ } // 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} @@ -665,10 +683,10 @@ do{ // Loop over the remaining incoming/waited messages 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 @@ -683,34 +701,34 @@ convergence. 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. @@ -719,30 +737,30 @@ Finally, the required modifications of the previous scheme 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) -double asyncStart; // Starting time of the current async section -int nbSyncIter = 0; // Number of sync iterations done in async mode +OpMode curMode = SYNC;// Current operating mode (always begin in sync) +double asyncStart; // Starting time of the current async section +int nbSyncIter = 0; // Number of sync iterations done in async mode // Computation loop while(!Finished){ @@ -751,14 +769,15 @@ while(!Finished){ // Entering synchronous mode when asyncDuration is reached @% // (additional conditions can be specified if needed) @ if(MPI_Wtime() - asyncStart >= asyncDuration){ - // Waiting for the end of previous sends before starting sync mode + // Waiting for the end of previous sends before starting sync mode omp_set_lock(&lockSendsDone); curMode = SYNC; // Entering synchronous mode stampData(dataToSend, SYNC); // Mark data to send with sync flag nbSyncIter = 0; } }else{ - // In main async mode, going back to async mode when the max number of sync iterations are done + // In main async mode, going back to async mode when the max number + // of sync iterations are done if(mainMode == ASYNC){ nbSyncIter++; // Update of the number of sync iterations done if(nbSyncIter == 2){ @@ -769,7 +788,7 @@ while(!Finished){ } } - // Sendings of data dependencies + // Sending of data dependencies if(curMode == SYNC || !SendsInProgress){ ... } @@ -780,18 +799,18 @@ while(!Finished){ } // 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 ... } } @@ -803,36 +822,38 @@ while(!Finished){ %\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 + // Determination of the dynamic operating mode, sendings of data + // 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 + ... // There may be several data transfers: typically A and b in + // 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; @@ -858,44 +879,44 @@ while(!Finished){ %\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 @@ -903,12 +924,13 @@ omp_lock_t lockRes; // Informs aux thread about new results omp_lock_t lockWrite; // Controls exclusion of results access ... auxRes ... ; // Results of auxiliary computations -// Parameters reading, MPI initialization, data initialization and distribution +// Parameters reading, MPI initialization, and data initialization and +// distribution ... // OpenMP initialization ... omp_init_lock(&lockAux); -omp_set_lock(&lockAux); // Unlocked when new aux results are available +omp_set_lock(&lockAux);//Unlocked when new aux results are available omp_init_lock(&lockRes); omp_set_lock(&lockRes); // Unlocked when new results are available omp_init_lock(&lockWrite); @@ -955,30 +977,32 @@ MPI_Finalize(); %\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<<>>(... @\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); @@ -996,13 +1020,13 @@ while(!Finished){ ... // 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 ... } } @@ -1013,8 +1037,8 @@ while(!Finished){ %\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 @@ -1027,13 +1051,13 @@ while(!Finished){ for(ind=0; ind