\section{Introduction}\label{ch6:intro}
-This chapter proposes to draw several development methodologies to obtain
+This chapter proposes to draw upon several development methodologies to obtain
efficient codes in classical scientific applications. Those methodologies are
-based on the feedback from several research works involving GPUs, either alone
-in a single machine or in a cluster of machines. Indeed, our past
-collaborations with industries have allowed us to point out that in their
-economical context, they can adopt a parallel technology only if its
-implementation and maintenance costs are small according to the potential
-benefits (performance, accuracy,...). So, in such contexts, GPU programming is
-still regarded with some distance according to its specific field of
-applicability (SIMD/SIMT model) and its still higher programming complexity and
-maintenance. In the academic domain, things are a bit different but studies for
-efficiently integrating GPU computations in multi-core clusters with maximal
-overlapping of computations with communications and/or other computations, are
-still rare.
+based on the feedback from several research works involving GPUs, either in a
+single machine or in a cluster of machines. Indeed, our past collaborations
+with industries have allowed us to point out that in their economical context,
+they can adopt a parallel technology only if its implementation and maintenance
+costs are small compared with the potential benefits (performance,
+accuracy, etc.). So, in such contexts, GPU programming is still regarded with
+some distance due to its specific field of applicability (SIMD/SIMT model:
+Single Instruction Multiple Data/Thread) and its still higher programming
+complexity and maintenance. In the academic domain, things are a bit different,
+but studies for efficiently integrating GPU computations in multicore clusters
+with maximal overlapping of computations with communications and/or other
+computations are still rare.
-For these reasons, the major aim of that chapter is to propose as simple as
-possible general programming patterns that can be followed or adapted in
+For these reasons, the major aim of that chapter is to propose general
+programming patterns, as simple as possible, that can be followed or adapted in
practical implementations of parallel scientific applications.
% Also, according to our experience in industrial collaborations, we propose a
% small prospect analysis about the perenity of such accelerators in the
% middle/long term.
-Also, we propose in a third part, a prospect analysis together with a particular
-programming tool that is intended to ease multi-core GPU cluster programming.
+In addition, we propose a prospect analysis together with a particular
+programming tool that is intended to ease multicore GPU cluster programming.
%%% Local Variables:
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
+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
+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}
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 non convergence (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 received 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.}
+%\pagebreak
+\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}
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.
% \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}
+\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
% \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}
+\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
...
}
}
% \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}
+\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}.
+non overlapped 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
-\clearpage
-\section{Perspective: A unifying programming model}
+%\clearpage
+\section{Perspective: a unifying programming model}
\label{sec:ch6p3unify}
In the previous sections we have seen that controlling a distributed GPU
-application when using tools that are commonly available is a quite challenging
+application when using tools that are commonly available is quite a challenging
task. To summarize, such an application has components that can be roughly
-classified as:
+classified as
\begin{description}
\itemsep=0pt
-\item[CPU:] CPU bound computations, realized as procedures in the chosen
+\item[CPU:] CPU-bound computations, realized as procedures in the chosen
programming language
-\item[CUDA$_{kern}$:] GPU bound computations, in our context realized as CUDA compute kernels
+\item[CUDA$_{kern}$:] GPU-bound computations, in our context realized as CUDA compute kernels
\item[CUDA$_{trans}$:] data transfer between CPU and GPU, realized with CUDA function calls
\item[MPI:] distributed data transfer routines, realized with MPI communication primitives
\item[OpenMP:] inter-thread control, realized with OpenMP synchronization tools such as mutexes
-\item[CUDA$_{sync}$] synchronization of the GPU, realized with CUDA functions
+\item[CUDA$_{sync}$:] synchronization of the GPU, realized with CUDA functions
\end{description}
Among these, the last (CUDA$_{sync}$) is not strictly necessary on modern
-systems, but still recommended to obtain optimal performance. With or without
+systems, but it is still recommended to obtain optimal performance. With or without
that last step, such an application is highly complex: it is difficult to design
or to maintain, and depends on a lot of different software components. The goal
-of this section is to present a new path of development that allows to replace
-the last three or four types of components that control the application (MPI,
-OpenMP, CUDA$_{sync}$ and eventually CUDA$_{trans}$) by a single tool:
-\textbf{O}rdered \textbf{R}ead-\textbf{W}rite \textbf{L}ocks, ORWL\index{ORWL},
-see~\cite{clauss:2010:inria-00330024:1,gustedt:hal-00639289}. Besides the
-simplification of the algorithmic scheme that we already have mentioned, the
-ongoing implementation of ORWL allows to use a feature of modern platforms that
-can improve the performance of CPU bound computations: lock-free atomic
+of this section is to present a new path of development that allows the
+replacement of the last three or four types of components that control the application (MPI,
+OpenMP, CUDA$_{sync}$, and eventually CUDA$_{trans}$) with a single tool:
+{O}rdered {R}ead-{W}rite {L}ocks, ORWL\index{ORWL}
+(see~\cite{clauss:2010:inria-00330024:1,gustedt:hal-00639289}). Besides the
+simplification of the algorithmic scheme that we have already mentioned, the
+ongoing implementation of ORWL allows the use of a feature of modern platforms that
+can improve the performance of CPU-bound computations: lock-free atomic
operations to update shared data consistently. For these, ORWL relies on new
interfaces that are available with the latest revision of the ISO standard for
-the C programming language, see~\cite{C11}.
+the C programming language (see~\cite{C11}).
\subsection{Resources}
\label{sec:ch6p3resources}
abstraction: \emph{resources}\index{ORWL!resource}. An ORWL resource may correspond to a local or
remote entity and is identified through a \emph{location}\index{ORWL!location}, that is a unique
identification through which it can be accessed from all different components of
-the same application. In fact, resources and locations (entities and their names
+the same application. In fact, resources and locations (entities and their names,
so to speak) are mostly identified by ORWL and these words will be used
interchangeably.
-Resources may be of very different kind:
+Resources may be of very different kinds:
\begin{description}
\item[Data] resources\index{ORWL!resource!data} are entities that represents data and not specific
- memory buffers or locations. During the execution of an application it can be
+ memory buffers or locations. During the execution of an application they can be
\emph{mapped} repeatedly into the address space and in effect be represented
at different addresses. Data resources can be made accessible uniformly in all
- parts of the application, provided that the locking protocol is observed, see
- below. Data resources can have different persistence:
+ parts of the application, provided that the locking protocol is observed (see
+ below). Data resources can have different persistence:
\begin{description}
\item[RAM] data resources are typically temporary data that serve only during
a single run of the application. They must be initialized at the beginning of
- their lifetime and the contents is lost at the end.
+ their lifetime and the contents are lost at the end.
\item[File] data resources are persistent and linked to a file in the file
system of the platform.
\item[Collective] data resources are data to which all tasks of an
application contribute (see below). Examples for such resources are \emph{broadcast},
- \emph{gather} or
- \emph{reduce} resources, e.g to distribute initial data or to collect the
+ \emph{gather}, or
+ \emph{reduce} resources, e.g., to distribute initial data or to collect the
result of a distributed computation.
\end{description}
- Other types of data resources could be easily implemented with ORWL, e.g web
+ Other types of data resources could be easily implemented with ORWL, e.g., web
resources (through a ftp, http or whatever server address) or fixed hardware
addresses.
\item[Device] resources\index{ORWL!resource!device} represent hardware entities of the
platform. ORWL can then be used to regulate the access to such device
resources. In our context the most important such resource is the GPU, but we
- could easily use it to represent a CPU core, a camera or other peripheral
+ could easily use it to represent a CPU core, a camera, or another peripheral
device.
\end{description}
\Lst{algo:ch6p3ORWLresources} shows an example of a declaration of
four resources per task. Two (\texttt{curBlock} and \texttt{nextBlock}) are
intended to represent the data in a block-cyclic parallel matrix multiplication
-(see p.~\pageref{ch6:p1block-cyclic}),
-\texttt{GPU} represents a GPU device and \texttt{result} will represent a
+(see Section~\ref{ch6:p1block-cyclic}),
+\texttt{GPU} represents a GPU device, and \texttt{result} will represent a
collective ``gather'' resource among all the tasks.
+\pagebreak
%\begin{algorithm}[H]
% \caption{Declaration of ORWL resources for a block-cyclic matrix
% multiplication.}
% \label{algo:ch6p3ORWLresources}
-\begin{Listing}{algo:ch6p3ORWLresources}{Declaration of ORWL resources for a block-cyclic matrix multiplication}
+\begin{Listing}{algo:ch6p3ORWLresources}{declaration of ORWL resources for a block-cyclic matrix multiplication}
#include "orwl.h"
...
ORWL_LOCATIONS_PER_TASK(curBlock, nextBlock, GPU, result);
ORWL_DATA_LOCATION(curBlock);
-ORWL_DATA_LOCATION(nexBlock);
+ORWL_DATA_LOCATION(nextBlock);
ORWL_DEVICE_LOCATION(GPU);
ORWL_GATHER_LOCATION(result);
\end{Listing}
\label{sec:ch6p3ORWLcontrol}
-ORWL regulates access to all its resources, no ``random access'' to a resource
+ORWL regulates access to all its resources; no ``random access'' to a resource
is possible. It doesn't even have a user-visible data type for resources.
\begin{itemize}
\item All access is provided through \emph{handle}s\index{ORWL!handle}. Similar to pointers or
links, these only refer to a resource and help to manipulate it. Usually
several handles to the same resource exist, even inside the same OS process
or thread, or in the same application task.
-\item The access is locked with RW semantics, where \emph{R} stands for
- concurrent \textbf{R}ead~access, and \emph{W} for exclusive
- \textbf{W}rite~access. This feature replaces the control aspect of MPI
- communications, OpenMP inter-thread control and CUDA$_{sync}$.
-\item This access is \textbf{O}rdered (or serialized)\index{ordered access} through a FIFO, \emph{one
+\item The access is locked with RW semantics, where {R} stands for
+ concurrent {R}ead~access, and {W} for exclusive
+ {W}rite~access. This feature replaces the control aspect of MPI
+ communications, OpenMP inter-thread control, and CUDA$_{sync}$.
+\item This access is {O}rdered (or serialized)\index{ordered access} through a FIFO, \emph{one
FIFO per resource}. This helps to run the different tasks of an application
in a controlled order and to always have all resources in a known state. This
aspect largely replaces and extends the ordering of tasks that MPI typically
\subsection{Example: block-cyclic matrix multiplication (MM)}
\label{sec:ch6p3ORWLMM}
-Let us now have a look how a block-cyclic matrix multiplication\index{matrix
+Let us now have a look at how a block-cyclic matrix multiplication\index{matrix
multiplication!block cyclic} algorithm can be
-implemented with these concepts (\Lst{algo:ch6p3ORWLBCCMM}). Inside the loop it mainly consists of three
-different operations, of which the first two can be run concurrently, and the
+implemented with these concepts (\Lst{algo:ch6p3ORWLBCCMM}). Inside the loop
+there are mainly three
+different operations, the first two of which can be run concurrently, and the
third must be done after the other two.
%\begin{algorithm}[H]
% \caption{Block-cyclic matrix multiplication, high level per task view.}
% \label{algo:ch6p3ORWLBCCMM}
-\begin{Listing}{algo:ch6p3ORWLBCCMM}{Block-cyclic matrix multiplication, high level per task view}
+\begin{Listing}{algo:ch6p3ORWLBCCMM}{block-cyclic matrix multiplication, high level per task view}
typedef double MBlock[N][N];
MBlock A;
MBlock B[k];
%\end{algorithm}
-\Lst{algo:ch6p3ORWLlcopy} shows the local copy operation~3 as it could
+\Lst{algo:ch6p3ORWLlcopy} shows the local copy operation~3 from line 17 of
+\Lst{algo:ch6p3ORWLBCCMM} as it could
be realized with ORWL.
It uses two resource handles
\texttt{nextRead} and \texttt{curWrite} and marks nested \emph{critical
sections} for these handles. Inside the nested sections it obtains pointers to
the resource data; the resource is \emph{mapped} into the address space of the
program, and then a standard call to \texttt{memcpy} achieves the operation
-itself. The operation is integrated in its own \textbf{for}-loop, such that it
+itself. The operation is integrated in its own {for}-loop, such that it
could run independently in an OS thread by its own.
%\begin{algorithm}[H]
% \caption{An iterative local copy operation.}
% \label{algo:ch6p3ORWLlcopy}
-\begin{Listing}{algo:ch6p3ORWLlcopy}{An iterative local copy operation}
+\begin{Listing}{algo:ch6p3ORWLlcopy}{an iterative local copy operation}
for (size_t i = 0; i < k; ++i) {
ORWL_SECTION(nextRead) {
MBlock const* sBlock = orwl_read_map(nextRead);
%\end{algorithm}
Next, in \Lst{algo:ch6p3ORWLrcopy} we copy data from a remote task to
-a local task. Substantially the operation is the same, only that in the example
-different handles (\texttt{remRead} and \texttt{nextWrite}) are used to
+a local task. Substantially the operation is the same, save for the
+different handles (\texttt{remRead} and \texttt{nextWrite}) that are used to
represent the respective resources.
%\begin{algorithm}[H]
% \caption{An iterative remote copy operation as part of a block cyclic matrix
% multiplication task.}
%\label{algo:ch6p3ORWLrcopy}
-\begin{Listing}{algo:ch6p3ORWLrcopy}{An iterative remote copy operation as part of a block cyclic matrix multiplication task}
+\begin{Listing}{algo:ch6p3ORWLrcopy}{an iterative remote copy operation as part of a block cyclic matrix multiplication task}
for (size_t i = 0; i < k; ++i) {
ORWL_SECTION(remRead) {
MBlock const* sBlock = orwl_read_map(remRead);
Now let us have a look into the operation that probably interests us the most,
the interaction with the GPU in \Lst{algo:ch6p3ORWLtrans}. Again there
-is much structural resemblance to the copy operations from above, only that we
+is much structural resemblance to the copy operations from above, but we
transfer the data to the GPU in the innermost block and then run the GPU MM
-kernel while we still are inside the critical section for the GPU.
+kernel while we are still inside the critical section for the GPU.
%\begin{algorithm}[H]
% \caption{An iterative GPU transfer and compute operation as part of a block cyclic matrix
% multiplication task.}
% \label{algo:ch6p3ORWLtrans}
-\begin{Listing}{algo:ch6p3ORWLtrans}{An iterative GPU transfer and compute operation as part of a block cyclic matrix multiplication task}
+\begin{Listing}{algo:ch6p3ORWLtrans}{an iterative GPU transfer and compute operation as part of a block cyclic matrix multiplication task}
for (size_t i = 0; i < k; ++i) {
ORWL_SECTION(GPUhandle) {
ORWL_SECTION(curRead) {
%\end{algorithm}
Now that we have seen how the actual procedural access to the resources is
-regulated we will show how the association between handles and resources is
-specified. E.g in our application of block-cyclic MM the \texttt{curRead} handle
-should correspond to current matrix block of the corresponding task, whereas
+regulated, we will show how the association between handles and resources is
+specified. In our application of block-cyclic MM the \texttt{curRead} handle
+should correspond to the current matrix block of the corresponding task, whereas
\texttt{remRead} should point to the current block of the neighboring task.
-Both read operations on these matrix blocks can be effected without creating
+Both read operations on these matrix blocks can be performed without creating
conflicts, so we would like to express that fact in our resource specification.
From a point of view of the resource ``current block'' of a particular task,
this means that it can have two simultaneous readers, the task itself performing
the transfer to the GPU, and the neighboring task transferring the data to its
-``next block''.
+``next block.''
\Lst{algo:ch6p3ORWLdecl} first shows the local dynamic declarations of
our application; it declares a \texttt{block} type for the matrix blocks, a
%\begin{algorithm}[H]
% \caption{Dynamic declaration of handles to represent the resources.}
% \label{algo:ch6p3ORWLdecl}
-\begin{Listing}{algo:ch6p3ORWLdecl}{Dynamic declaration of handles to represent the resources}
+\begin{Listing}{algo:ch6p3ORWLdecl}{dynamic declaration of handles to represent the resources}
/* A type for the matrix blocks */
typedef double MBlock[N][N];
/* Declaration to handle the collective resource */
\end{Listing}
%\end{algorithm}
-With these declarations, we didn't yet tell ORWL much about the resources to
-which these handles refer, nor the type (read or write) or the priority (FIFO
-position) of the access. This is done in code
-\Lst{algo:ch6p3ORWLinit}. The handles for
-\Lst{algo:ch6p3ORWLtrans} are given first, \texttt{GPUhandle} will be
-accessed exclusively (therefore the \texttt{write}) and, as said,
-\texttt{curRead} is used shared (so a \texttt{read}). Both are inserted in the
-FIFO of there respective resources with highest priority, specified by the
-\texttt{0}s in the third function parameter. The resources to which they
-correspond are specified through calls to the macro \texttt{ORWL\_LOCATION},
-indicating the task (\texttt{orwl\_mytid} is the ID of the current task) and the
-specific resource of that task, here \texttt{GPU} and \texttt{curBlock}.
-
-Likewise, a second block of insertions concerns the handles for
-\Lst{algo:ch6p3ORWLrcopy}: \texttt{newWrite} reclaims an exclusive
-access and \texttt{remRead} a shared. \texttt{remRead} corresponds to a
-resource of another task; \texttt{previous(orwl\_mytid)} is supposed to return
-the ID of the previous task in the cycle. Both accesses can be effected
-concurrently with the previous operation, so we insert with the same priority
-\texttt{0} as before.
-
-Then, for the specification of the third operation
-(\Lst{algo:ch6p3ORWLlcopy}) we need to use a different priority: the
-copy operation from \texttt{nextBlock} to \texttt{curBlock} has to be performed
-after the other operations have terminated.
-
-As a final step, we then tell ORWL that the specification of all accesses is
+With these declarations, we didn't yet tell ORWL much about the resources to
+which these handles refer, nor the type (read or write) or the priority (FIFO
+position) of the access. This is done in code \Lst{algo:ch6p3ORWLinit} that
+shows six insertions of handles into their respective FIFO locations. The
+handles \texttt{GPUhandle} and \texttt{curRead} are given first (lines 3 and 4),
+\texttt{GPUhandle} will be accessed exclusively (therefore the \texttt{write})
+and, as said, \texttt{curRead} is used in shared access (so a
+\texttt{read}). Both are inserted in the FIFO of their respective resources with
+highest priority, specified by the \texttt{0}s in the third function
+parameter. The resources to which they correspond are specified through calls to
+the macro \texttt{ORWL\_LOCATION}, indicating the task (\texttt{orwl\_mytid} is
+the ID of the current task) and the specific resource of that task, here
+\texttt{GPU} and \texttt{curBlock}.
+
+Likewise, a second block of insertions concerns the handles \texttt{remRead} and
+\texttt{nextWrite} (lines 8 and 9). \texttt{nextWrite} reclaims an exclusive
+access and \texttt{remRead} a shared one. \texttt{remRead} corresponds to a
+resource of another task; the call to \texttt{previous(orwl\_mytid)} is supposed
+to return the ID of the previous task in the cycle. Both accesses can be
+performed concurrently with the previous operation, so we insert them with the
+same priority \texttt{0} as previously.
+
+Then, for the specification of the third operation related to handles
+\texttt{nextRead} and \texttt{curWrite} (lines 13 and 14), we need to
+use a different priority: the copy operation from data locations \texttt{nextBlock} to
+\texttt{curBlock} has to be performed after the other operations have
+terminated (so the priority \texttt{1}).
+
+As a final step, we then tell ORWL (line 16) that the specification of all accesses is
complete and that it may schedule\index{ORWL!schedule} all these accesses in the respective FIFOs of
the resources.
%\begin{algorithm}[H]
% \caption{Dynamic initialization of access mode and priorities.}
% \label{algo:ch6p3ORWLinit}
-\pagebreak
-\begin{Listing}{algo:ch6p3ORWLinit}{Dynamic initialization of access mode and priorities}
+%\pagebreak
+\begin{Listing}{algo:ch6p3ORWLinit}{dynamic initialization of access mode and priorities}
/* One operation with priority 0 (highest) consists */
-/* in copying from current to the GPU and run MM, there. */
+/* of copying from current to the GPU and running MM, there. */
orwl_write_insert(&GPUhandle, ORWL_LOCATION(orwl_mytid, GPU), 0);
orwl_read_insert(&curRead, ORWL_LOCATION(orwl_mytid, curBlock), 0);
/* Another operation with priority 0 consists */
-/* in copying from remote to next */
+/* of copying from remote to next */
orwl_read_insert(&remRead, ORWL_LOCATION(previous(orwl_mytid), curBlock), 0);
-orwl_write_insert(&nextWrite, ORWL_LOCATION(orwl_mytid, nextBlock), 0);
+orwl_write_insert(&nextWrite, ORWL_LOCATION(orwl_mytid,nextBlock), 0);
/* One operation with priority 1 consists */
-/* in copying from next to current */
+/* of copying from next to current */
orwl_read_insert(&nextRead, ORWL_LOCATION(orwl_mytid, nextBlock), 1);
orwl_write_insert(&curWrite, ORWL_LOCATION(orwl_mytid, curBlock), 1);
\subsection{Tasks and operations}
\label{sec:ch6p3tasks}
-With that example we have now seen that ORWL distinguishes
+With this example we have now seen that ORWL distinguishes
\emph{tasks}\index{ORWL!task} and
\emph{operations}\index{ORWL!operation}. An ORWL program is divided into tasks that can be seen as the
algorithmic units that will concurrently access the resources that the program
uses. A task for ORWL is characterized by
\begin{itemize}
-\item a fixed set of resources that it manages, ``\emph{owns}'', in our example
+\item a fixed set of resources that it manages, ``\emph{owns},'' in our example
the four resources that are declared in \Lst{algo:ch6p3ORWLresources}.
\item a larger set of resources that it accesses, in our example all resources
that are used in \Lst{algo:ch6p3ORWLinit}.
\item a set of operations that act on these resources, in our example the three
operations that are used in \Lst{algo:ch6p3ORWLBCCMM}, and that are
- elaborated in Listings~\ref{algo:ch6p3ORWLlcopy},~\ref{algo:ch6p3ORWLrcopy}
+ elaborated in Listings~\ref{algo:ch6p3ORWLlcopy},~\ref{algo:ch6p3ORWLrcopy},
and~\ref{algo:ch6p3ORWLtrans}.
\end{itemize}
\noindent {\bf Considered parallel algorithms and implementations}
\medskip
-This section focusses on synchronous parallel algorithms implemented with
+This section focuses on synchronous parallel algorithms implemented with
overlapping computations and communications\index{overlap!computation and communication}. Parallel synchronous algorithms are
-easier to implement, debug and maintain than asynchronous ones, see
-Section~\ref{ch6:part2}. Usually, they follow a BSP-like parallel
-scheme\index{BSP parallel scheme},
-alternating local computation steps and communication steps,
-see~\cite{Valiant:BSP}. Their execution is usually deterministic, excepted
+easier to implement, debug, and maintain than asynchronous ones (see
+Section~\ref{ch6:part2}). Usually, they follow a BSP-like parallel
+scheme\index{BSP parallel scheme} (Bulk Synchronous Parallel model),
+alternating local computation steps and communication steps
+(see~\cite{Valiant:BSP}). Their execution is usually deterministic, except
for stochastic algorithms that contain random number generations. Even in
-this case, their execution can be controlled during debug steps, allowing to
-track and to fix bugs quickly.
+this case, their execution can be controlled during debug steps, allowing the
+user to track and to fix bugs quickly.
However, depending on the properties of the algorithm, it is sometimes possible to
overlap computations and communications. If processes exchange data
that is not needed for the
-computation that is following immediately, it is possible to implement such
+computation immediately following, it is possible to implement such
an overlap. We have investigated the efficiency of this approach in previous
works~\cite{GUSTEDT:2007:HAL-00280094:1,ChVCV13}, using standard parallel
programming tools to achieve the implementation.
-The normalized and well known Message Passing Interface (MPI\index{MPI}) includes some asynchronous
-point-to-point communication routines, that should allow to implement some
+The normalized and well-known Message Passing Interface (MPI\index{MPI}) includes some asynchronous
+point-to-point communication routines, that should allow the implementation of some
communication/computation overlap. However, current MPI implementations do not achieve
-that goal efficiently; effective overlapping with MPI requires a group of
+this goal efficiently; effective overlapping with MPI requires a group of
dedicated threads (in our case implemented with OpenMP\index{OpenMP}) for the basic
synchronous communications while another group of threads executes computations
in parallel.
% decrease the execution time.
Nevertheless, communication and computation are not completely independent on
modern multicore architectures: they use shared hardware components such as the
-interconnection bus and the RAM. Therefore that approach only saved up to $20\%$
+interconnection bus and the RAM. Therefore, this approach saves only up to $20\%$
of the expected time on such a platform. This picture changes on clusters
-equipped with GPU. They effectively allow independence of computations on the
-GPU and communication on the mainboard (CPU, interconnection bus, RAM, network
-card). We saved up to $100\%$ of the expected time on our GPU cluster, as we
+equipped with GPUs. Indeed, GPUs effectively allow independence of computations they
+perform and communications done on the mainboard (CPU, interconnection bus, RAM, network
+card). We save up to $100\%$ of the expected time on our GPU cluster, as we
will expose in the next section.
\noindent {\bf Specific interests in GPU clusters}
\medskip
-In a computing node, a GPU is a kind of scientific coprocessor usually located
-on an auxiliary board, with its own memory. So, when data have been transferred
-from the CPU memory to the GPU memory, then GPU computations can be achieved on
-the GPU board, totally in parallel of any CPU activities (like internode cluster
+In a computing node, a GPU is a kind of scientific coprocessor, usually located
+on an auxiliary board, with its own memory. So, once data are transferred
+from the CPU memory to the GPU memory, GPU computations can be achieved on
+the GPU board, totally in parallel with any CPU activities (such as internode cluster
communications). The CPU and the GPU access their respective memories and do not
-interfere, so they can achieve a very good overlap\index{overlap!computation
-and computation} of their activities (better
-than two CPU cores).
+interfere with each other, so they can achieve a very good overlap\index{overlap!computation
+and computation} of their activities (better than two CPU cores).
-But using a GPU on a computing node requires to transfer data from the CPU to
-the GPU memory, and to transfer the computation results back from the GPU to the
+But using a GPU on a computing node requires the transfer of data from the CPU to
+the GPU memory, as well as the transfer of the computation results back from the GPU to the
CPU. Transfer times are not excessive, but depending on the application
they still can be significant compared to the GPU computation times. So, sometimes it
can be interesting to overlap the internode cluster communications with both the
parallel programming schemes on a GPU cluster:
\begin{enumerate}
-\item parallelizing only 'internode CPU communications' with 'GPU computations',
+\item parallelizing only internode CPU communications with GPU computations,
and achieving CPU/GPU data transfers before and after this parallel step,
-\item parallelizing 'internode CPU communications' with a '(sequential) sequence
- of CPU/GPU data transfers and GPU computations',
-\item parallelizing 'internode CPU communications' with a 'streamed sequence of
- CPU/GPU data transfers and GPU computations',
-\item parallelizing 'internode CPU communications' with 'CPU/GPU data transfers'
- and with 'GPU computations', interleaving computation-communication
+\item parallelizing internode CPU communications with a (sequential) sequence
+ of CPU/GPU data transfers and GPU computations,
+\item parallelizing internode CPU communications with a streamed sequence of
+ CPU/GPU data transfers and GPU computations,
+\item parallelizing internode CPU communications with CPU/GPU data transfers
+ and with GPU computations, interleaving computation-communication
iterations.
\end{enumerate}
\label{fig:ch6p1overlapnative}
\end{figure}
-Using CUDA\index{CUDA}, GPU kernel executions are non-blocking, and GPU/CPU data
+Using CUDA\index{CUDA}, GPU kernel executions are nonblocking, and GPU/CPU data
transfers\index{CUDA!data transfer}
-are blocking or non-blocking operations. All GPU kernel executions and CPU/GPU
-data transfers are associated to "streams"\index{CUDA!stream}, and all operations on a same stream
+are blocking or nonblocking operations. All GPU kernel executions and CPU/GPU
+data transfers are associated to "streams,"\index{CUDA!stream} and all operations on a same stream
are serialized. When transferring data from the CPU to the GPU, then running GPU
-computations and finally transferring results from the GPU to the CPU, there is
+computations, and finally transferring results from the GPU to the CPU, there is
a natural synchronization and serialization if these operations are achieved on
the same stream. GPU developers can choose to use one (default) or several
streams. In this first scheme of overlapping, we consider parallel codes using
only one GPU stream.
-"Non-blocking GPU kernel execution" means a CPU routine runs a parallel
+``Nonblocking GPU kernel execution'' means a CPU routine runs a parallel
execution of a GPU computing kernel, and the CPU routine continues its execution
(on the CPU) while the GPU kernel is running (on the GPU). Then the CPU routine
-can initiate some communications with some others CPU, and so it automatically
+can initiate some communications with some other CPUs, and so it automatically
overlaps the internode CPU communications with the GPU computations (see
\Fig{fig:ch6p1overlapnative}). This overlapping is natural when programming with
CUDA and MPI: it is easy to deploy, but does not overlap the CPU/GPU data
transfers.
+\Lst{algo:ch6p1overlapnative} introduces the generic code of a MPI+CUDA
+implementation, natively and implicitly overlapping MPI communications with CUDA
+GPU computations.
+
+
%\begin{algorithm}[t]
% \caption{Generic scheme implicitly overlapping MPI communications with CUDA GPU
% computations}\label{algo:ch6p1overlapnative}
-\pagebreak
-\begin{Listing}{algo:ch6p1overlapnative}{Generic scheme implicitly overlapping MPI communications with CUDA GPU computations}
+%\pagebreak
+\begin{Listing}{algo:ch6p1overlapnative}{generic scheme implicitly overlapping MPI communications with CUDA GPU computations}
// Input data and result variables and arrays (example with
// float datatype, 1D input arrays, and scalar results)
float *cpuInputTabAdr, *gpuInputTabAdr;
cpuResTabAdr = malloc(sizeof(float)*NbIter);
cudaMalloc(&gpuResAdr,sizeof(float));
-// Definition of the Grid of blocks of GPU threads
+// Definition of the grid of blocks of GPU threads
dim3 Dg, Db;
Dg.x = ...
...
\end{Listing}
%\end{algorithm}
-\Lst{algo:ch6p1overlapnative} introduces the generic code of a MPI+CUDA
-implementation, natively and implicitly overlapping MPI communications with CUDA
-GPU computations. Some input data and output results arrays and variables are
-declared and allocated from line 1 up to 10, and a computation loop is
-implemented from line 21 up to 34. At each iteration:
+
+
+Some input data and output results arrays and variables are
+declared and allocated from line~3 through 10, and a computation loop is
+implemented from line~22 through 34. At each iteration:
\begin{itemize}
-\item \texttt{cudaMemcpy} on line 23 transfers data from the CPU memory
+\item \texttt{cudaMemcpy} on line~23 transfers data from the CPU memory
to the GPU memory. This is a basic and synchronous data transfer.
-\item \texttt{gpuKernel\_k1<<<Dg,Db>>>} on line 26 starts GPU computation
- (running a GPU kernel on the grid of blocks of threads defined at line 12 to
- 15). This is a standard GPU kernel run, it is an asynchronous operation. The
+\item \texttt{gpuKernel\_k1<<<Dg,Db>>>} on line~26 starts GPU computation
+ (running a GPU kernel on the grid of blocks of threads defined in lines~13 to
+ 15). This is a standard GPU kernel run; it is an asynchronous operation. The
CPU can continue to run its code.
-\item \texttt{MPI\_Sendrecv\_replace} on line 27 achieves some blocking
- internode communications, overlapping GPU computations started just before.
-\item If needed, \texttt{cudaMemcpy} on line 31 transfers the iteration result from
- one variable in the GPU memory at one array index in the CPU memory (in this example the CPU
+\item \texttt{MPI\_Sendrecv\_replace} on line~27 achieves some blocking
+ internode communications, overlapping GPU computations started just previously.
+\item If needed, \texttt{cudaMemcpy} on line~31 transfers the iteration result from
+ one variable in the GPU memory to one array index in the CPU memory (in this example the CPU
collects all iteration results in an array). This operation is started after
the end of the MPI communication (previous instruction) and after the end of
the GPU kernel execution. CUDA insures an implicit synchronization of all
operations involving the same GPU stream, like the default stream in this
- example. Result transfer has to wait the GPU kernel execution is finished.
- If there is no result transfer implemented, the next operation on the GPU
- will wait until the GPU kernel execution will be ended.
+ example. The transfer of the results has to wait until the GPU kernel execution is finished.
+ If there is no results transfer implemented, the next operation on the GPU
+ will wait until the GPU kernel execution has ended.
\end{itemize}
This implementation is the easiest one involving the GPU. It achieves an
-implicit overlap of internode communications and GPU computations, no explicit
-multithreading is required on the CPU. However, CPU/GPU data transfers are
+implicit overlap of internode communications and GPU computations with no explicit
+multithreading required on the CPU. However, CPU/GPU data transfers are
achieved serially and not overlapped.
can be interesting to overlap internode CPU computations with a \emph{GPU
sequence}\index{GPU sequence} including CPU/GPU data transfers and GPU computations (see
\Fig{fig:ch6p1overlapseqsequence}). Algorithmic issues of this approach are basic,
-but their implementation require explicit CPU multithreading and
+but their implementation requires explicit CPU multithreading and
synchronization, and CPU data buffer duplication. We need to implement two
-threads, one starting and achieving MPI communications, and the other running
+threads, one starting and achieving MPI communications and the other running
the \emph{GPU sequence}. OpenMP allows an easy and portable implementation of
this overlapping strategy. However, it remains more complex to develop and to
maintain than the previous strategy (overlapping only internode CPU
-communications and GPU computations), and should be adopted only when CPU/GPU
+communications and GPU computations) and should be adopted only when CPU/GPU
data transfer times are not negligible.
+\Lst{algo:ch6p1overlapseqsequence} introduces the generic code of a
+MPI+OpenMP+CUDA implementation, explicitly overlapping MPI communications with
+GPU sequences.
+
%\begin{algorithm}
% \caption{Generic scheme explicitly overlapping MPI communications with
% sequences of CUDA CPU/GPU transfers and CUDA GPU
% computations}\label{algo:ch6p1overlapseqsequence}
-\begin{Listing}{algo:ch6p1overlapseqsequence}{Generic scheme explicitly overlapping MPI communications with sequences of CUDA CPU/GPU transfers and CUDA GPU computations}
+\begin{Listing}{algo:ch6p1overlapseqsequence}{generic scheme explicitly overlapping MPI communications with sequences of CUDA CPU/GPU transfers and CUDA GPU computations}
// Input data and result variables and arrays (example with
// float datatype, 1D input arrays, and scalar results)
float *cpuInputTabAdrCurrent, *cpuInputTabAdrFuture, *gpuInputTabAdr;
cpuResTabAdr = malloc(sizeof(float)*NbIter);
cudaMalloc(&gpuResAdr,sizeof(float));
-// Definition of the Grid of blocks of GPU threads
+// Definition of the grid of blocks of GPU threads
dim3 Dg, Db;
Dg.x = ...
...
cudaMemcpyDeviceToHost);
}
- // - Wait for both threads have achieved their iteration tasks
+ // - Wait until both threads have achieved their iteration tasks
#pragma omp barrier
- // - Each thread permute its local buffer pointers
+ // - Each thread permutes its local buffer pointers
tmp = current;
current = future;
future = tmp;
\end{Listing}
%\end{algorithm}
-\Lst{algo:ch6p1overlapseqsequence} introduces the generic code of a
-MPI+OpenMP+CUDA implementation, explicitly overlapping MPI communications with
-\emph{GPU sequences}. Lines 25--62 implement the OpenMP parallel region,
-around the computation loop (lines 33--61). For performances it is
+Lines~25--62 implement the OpenMP parallel region,
+around the computation loop (lines~33--61). For efficient performances it is
important to create and destroy threads only one time (not at each iteration):
-the parallel region has to surround the computation loop. Lines 1--11
-consist in declaration and allocation of input data arrays and result arrays and
-variables, like in previous algorithm (\Lst{algo:ch6p1overlapnative}). However, we implement two input data buffers on the
+the parallel region has to surround the computation loop. Lines~3--11
+consist of declaration and allocation of input data arrays and result arrays and
+variables, as in the previous algorithm (\Lst{algo:ch6p1overlapnative}). However, we implement two input data buffers on the
CPU (current and future version). As we aim to overlap internode MPI
communications and GPU sequence, including CPU to GPU data transfer of current
input data array, we need to store the received new input data array in a
separate buffer. Then, the current input data array will be safely read on the
CPU and copied into the GPU memory.
-The thread creations\index{OpenMP!thread creation} are easily achieved with one OpenMP directive (line
-25). Then each thread defines and initializes \emph{its} local buffer pointers,
-and enters \emph{its} computing loop (lines 27--33). Inside the computing
-loop, a test on the thread number allows to run a different code in each
-thread. Lines 37--40 implement the MPI synchronous communication run by
-thread number $0$. Lines 45--52 implement the GPU sequence run by thread
-$1$: CPU to GPU data transfer, GPU computation and GPU to CPU result
+The thread creations\index{OpenMP!thread creation} are easily achieved with one OpenMP directive (line~25).
+Then each thread defines and initializes \emph{its} local buffer pointers,
+and enters \emph{its} computing loop (lines~28--33). Inside the computing
+loop, a test on the thread number makes it possible to run a different code in each
+thread. Lines~37--40 implement the MPI synchronous communication run by
+thread number $0$. Lines~45--52 implement the GPU sequence run by thread
+$1$: CPU to GPU data transfer, GPU computation, and GPU to CPU result
transfer (if needed). Details of the three operations of this sequence have not changed
-compared to the previous overlapping strategy.
+from the previous overlapping strategy.
At the end of \Lst{algo:ch6p1overlapseqsequence}, an OpenMP synchronization
-barrier\index{OpenMP!barrier} on line 56 allows to wait OpenMP threads have achieved MPI
-communications and GPU sequence, and do not need to access the current input
-data buffer. Then, each thread permute its local buffer pointers (lines 58--60),
+barrier\index{OpenMP!barrier} on line~56 forces the OpenMP threads to wait until
+MPI
+communications and GPU sequence are achieved. %, and do not need to access the current input data buffer.
+Then, each thread permutes its local buffer pointers (lines~58--60),
and is ready to enter the next iteration, processing the new current input
array.
-
%\subsubsection{Overlapping with a streamed GPU sequence}
\bigskip
+%\pagebreak
\noindent {\bf Overlapping with a streamed GPU sequence}
\medskip
+Depending on the algorithm implemented, it is sometimes possible to split the
+GPU computation into several parts processing distinct data. Then, we can
+speedup the GPU sequence using several CUDA streams\index{CUDA!stream}. The goal is
+to overlap CPU/GPU data transfers with GPU computations\index{overlap!GPU data transfers with GPU computation} inside the GPU
+ sequence. Compared to the previous overlapping strategy, we have to split the
+initial data transfer into a set of $n$ asynchronous and smaller data transfers,
+and split the initial GPU kernel call into a set of $n$ calls to the same GPU
+kernel. Usually, these smaller calls are deployed with fewer GPU threads
+(i.e., associated to a smaller grid of blocks of threads). Then, the first GPU
+computations can start as soon as the first data transfer has been achieved, and
+next transfers can be done in parallel with next GPU computations (see
+\Fig{fig:ch6p1overlapstreamsequence}).
+
\begin{figure}[t]
\centering
\includegraphics[width=\columnwidth]{Chapters/chapter6/figures/Sync-StreamSequenceOverlap.pdf}
\label{fig:ch6p1overlapstreamsequence}
\end{figure}
-Depending on the algorithm implemented, it is sometimes possible to split the
-GPU computation into several parts processing distinct data. Then, we can
-speedup the \emph{GPU sequence} using several CUDA \emph{streams}\index{CUDA!stream}. The goal is
-to overlap CPU/GPU data transfers with GPU computations\index{overlap!GPU data transfers with GPU computation} inside the \emph{GPU
- sequence}. Compared to the previous overlapping strategy, we have to split the
-initial data transfer in a set of $n$ asynchronous and smaller data transfers,
-and to split the initial GPU kernel call in a set of $n$ calls to the same GPU
-kernel. Usually, these smaller calls are deployed with less GPU threads
-(i.e. associated to a smaller grid of blocks of threads). Then, the first GPU
-computations can start as soon as the first data transfer has been achieved, and
-next transfers can be done in parallel of next GPU computations (see
-\Fig{fig:ch6p1overlapstreamsequence}).
-
-NVIDIA advises to start all asynchronous CUDA data transfers, and then to call
+NVIDIA advises starting all asynchronous CUDA data transfers and then calling
all CUDA kernel executions, using up to $16$ streams \cite{cudabestpractices}.
-Then, CUDA driver and
-runtime optimize the global execution of these operations. So, we cumulate two
-overlapping mechanisms. The former is controlled by CPU multithreading, and
-overlap MPI communications and the \emph{streamed GPU sequence}. The latter is
-controlled by CUDA programming, and overlap CPU/GPU data transfers and GPU
-computations. Again, OpenMP allows to easily implement the CPU multithreading,
-and to wait for the end of both CPU threads before to execute the next instructions
+Then, the CUDA driver and
+runtime optimize the global execution of these operations. So, we accumulate two
+overlapping mechanisms. The former is controlled by CPU multithreading and
+overlaps MPI communications and the streamed GPU sequence. The latter is
+controlled by CUDA programming and overlaps CPU/GPU data transfers and GPU
+computations. Again, OpenMP allows the easy implementation of the CPU multithreading
+and waiting for the end of both CPU threads before executing the next instructions
of the code.
+\Lst{algo:ch6p1overlapstreamsequence} introduces the generic MPI+OpenMP+CUDA
+code, explicitly overlapping MPI communications with
+streamed GPU sequences\index{GPU sequence!streamed}.
+
%\begin{algorithm}
% \caption{Generic scheme explicitly overlapping MPI communications with streamed sequences of CUDA
% CPU/GPU transfers and CUDA GPU computations}\label{algo:ch6p1overlapstreamsequence}
-\begin{Listing}{algo:ch6p1overlapstreamsequence}{Generic scheme explicitly overlapping MPI communications with streamed sequences of CUDA CPU/GPU transfers and CUDA GPU computations}
+\begin{Listing}{algo:ch6p1overlapstreamsequence}{generic scheme explicitly overlapping MPI communications with streamed sequences of CUDA CPU/GPU transfers and CUDA GPU computations}
// Input data and result variables and arrays (example with
// float datatype, 1D input arrays, and scalar results)
float *cpuInputTabAdrCurrent, *cpuInputTabAdrFuture, *gpuInputTabAdr;
cudaStream_t TabS[NbS];
for(int s = 0; s < NbS; s++)
cudaStreamCreate(&TabS[s]);
-// Definition of the Grid of blocks of GPU threads
+// Definition of the grid of blocks of GPU threads
...
// Set the number of OpenMP threads (to create) to 2
omp_set_num_threads(2);
sizeof(float),
cudaMemcpyDeviceToHost);
}
- // - Wait for both threads have achieved their iteration tasks
+ // - Wait until both threads have achieved their iteration tasks
#pragma omp barrier
- // - Each thread permute its local buffer pointers
+ // - Each thread permutes its local buffer pointers
tmp = current; current = future; future = tmp;
} // End of computation loop
} // End of OpenMP parallel region
\end{Listing}
%\end{algorithm}
-\Lst{algo:ch6p1overlapstreamsequence} introduces the generic MPI+OpenMP+CUDA
-code explicitly overlapping MPI communications with
-\emph{streamed GPU sequences}\index{GPU sequence!streamed}. Efficient usage of CUDA \emph{streams} requires to execute
-asynchronous CPU/GPU data transfers, that needs to read page-locked
+Efficient usage of CUDA streams requires executing
+asynchronous CPU/GPU data transfers, which implies reading page-locked
data\index{page-locked data} in CPU memory. So, CPU
-memory allocations on lines 6 and 7 are implemented with \texttt{cudaHostAlloc} instead of
-the basic \texttt{malloc} function. Then, $NbS$ \emph{streams} are created on lines 12--14.
+memory allocations on lines~6 and 7 are implemented with \texttt{cudaHostAlloc} instead of
+the basic \texttt{malloc} function. Then, $NbS$ \emph{streams} are created on lines~12--14.
Usually we create $16$ streams: the maximum number supported by CUDA.
-An OpenMP parallel region\index{OpenMP!parallel region} including two threads is implemented on lines 17--61 of
-\Lst{algo:ch6p1overlapstreamsequence}, similarly to the previous algorithm (see
+An OpenMP parallel region\index{OpenMP!parallel region} including two threads is implemented on lines~18--61 of
+\Lst{algo:ch6p1overlapstreamsequence}, as in the previous algorithm (see
\Lst{algo:ch6p1overlapseqsequence}). Code of thread $0$ achieving MPI communication is unchanged, but
-code of thread $1$ is now using streams. Following NVIDIA recommandations, we have first implemented
-a loop starting $NbS$ asynchronous data transfers (lines 39--45): transferring $N/NbS$ data on
-each stream. Then we have implemented a second loop (lines 46--48), starting asynchronous
+code of thread $1$ is now using streams. Following NVIDIA recommandations, we first implement
+a loop starting $NbS$ asynchronous data transfers (lines~39--45): transferring $N/NbS$ data on
+each stream. Then we implement a second loop (lines~46--48), starting asynchronous
executions of $NbS$ grids of blocks of GPU threads (one per stream). Data transfers and kernel
executions on the same stream are synchronized by CUDA and the GPU. So, each kernel execution will
-start after its data will be transferred into the GPU memory, and the GPU scheduler ensures to start
+start after its data has been transferred into the GPU memory, and the GPU
+scheduler ensures the start of
some kernel executions as soon as the first data transfers are achieved. Then, next data transfers
will be overlapped with GPU computations. After the kernel calls, on the different streams,
we wait for the end of all GPU threads previously run, calling an explicit synchronization
-function on line 49. This synchronization is not mandatory, but it will make the implementation more
+function on line~49. This synchronization is not mandatory, but it will make the implementation more
robust and will facilitate the debugging steps: all GPU computations run by the OpenMP thread number
-$1$ will be achieved before this thread will enter a new loop iteration, or before the computation
-loop will be ended.
+$1$ will be achieved before this thread enters a new loop iteration, or before the computation
+loop has ended.
If a partial result has to be transferred from GPU to CPU memory at the end of each loop iteration
-(for example the result of one \emph{reduction} per iteration), this transfer is achieved
-synchronously on the default stream (no particular stream is specified) on lines 51--54.
-Availability of the result values is ensured by the synchronization implemented on line 49.
-However, if a partial result has to be transferred on the CPU on each stream, then $NbS$ asynchronous data
-transfers could be started in parallel (one per stream), and should be implemented before the
-synchronization operation on line 49. The end of the computation loop includes a synchronization
-barrier of the two OpenMP threads, waiting they have finished to access the different data
-buffers in the current iteration. Then, each OpenMP thread exchanges its local buffer pointers, like
-in the previous algorithm. However, after the computation loop, we have added the
-destruction of the CUDA streams (lines 63--65).
-
-Finally, CUDA streams\index{CUDA!stream} have been used to extend \Lst{algo:ch6p1overlapseqsequence}
-with respect to its global scheme. \Lst{algo:ch6p1overlapstreamsequence} still creates an
-OpenMP parallel region, with two CPU threads, one in charge of MPI communications, and the other
-managing data transfers and GPU computations. Unfortunately, using GPU streams require to be able to
-split a GPU computation in independent subparts, working on independent subsets of data.
-\Lst{algo:ch6p1overlapstreamsequence} is not so generic than \Lst{algo:ch6p1overlapseqsequence}.
-
-
-\subsection{Interleaved communications-transfers-computations overlapping}
-
-\begin{figure}[t]
- \centering
- \includegraphics{Chapters/chapter6/figures/Sync-CompleteInterleaveOverlap.pdf}
- \caption{Complete overlap of internode CPU communications, CPU/GPU data transfers and GPU
- computations, interleaving computation-communication iterations}
- \label{fig:ch6p1overlapinterleaved}
-\end{figure}
-
-Many algorithms do not support to split data transfers and kernel calls, and can
-not exploit CUDA streams. For example, when each GPU thread requires to access
+(for example, the result of one \emph{reduction} per iteration), this transfer is achieved
+synchronously on the default stream (no particular stream is specified) on lines~51--54.
+Availability of the result values is ensured by the synchronization implemented on line~49.
+However, if a partial result has to be transferred onto the CPU on each stream, then $NbS$ asynchronous data
+transfers could be started in parallel (one per stream) and should be implemented before the
+synchronization operation on line~49. The end of the computation loop includes a synchronization
+barrier of the two OpenMP threads, waiting until they have finished accessing the different data
+buffers in the current iteration. Then, each OpenMP thread exchanges its local buffer pointers, as
+in the previous algorithm. After the computation loop, we have added the
+destruction of the CUDA streams (lines~64--65).
+
+In conclusion, CUDA streams\index{CUDA!stream} have been used to extend
+\Lst{algo:ch6p1overlapseqsequence} with respect to its global
+scheme. \Lst{algo:ch6p1overlapstreamsequence} still creates an OpenMP parallel
+region, with two CPU threads, one in charge of MPI communications and the other
+managing data transfers and GPU computations. Unfortunately, using GPU streams
+requires the ability to split a GPU computation into independent subparts,
+working on independent subsets of data. \Lst{algo:ch6p1overlapstreamsequence}
+is not so generic as \Lst{algo:ch6p1overlapseqsequence}.
+
+
+\subsection{Interleaved communications-transfers-computations\\overlapping}
+
+Many algorithms do not support splitting data transfers and kernel calls, and
+cannot exploit CUDA streams, for example, when each GPU thread requires access to
some data spread in the global set of transferred data. Then, it is possible to
-overlap internode CPU communications and CPU/GPU data transfers and GPU
+overlap internode CPU communications, CPU/GPU data transfers, and GPU
computations, if the algorithm achieves \emph{computation-communication
iterations} and if we can interleave these iterations. At iteration $k$: CPUs
exchange data $D_k$, each CPU/GPU couple transfers data $D_k$, and each GPU
strategy requires twice as many CPU data buffers
and twice as many GPU buffers.
+\begin{figure}[t]
+ \centering
+ \includegraphics{Chapters/chapter6/figures/Sync-CompleteInterleaveOverlap.pdf}
+ \caption{Complete overlap of internode CPU communications, CPU/GPU data transfers, and GPU
+ computations, interleaving computation-communication iterations.}
+ \label{fig:ch6p1overlapinterleaved}
+\end{figure}
+
+\Lst{algo:ch6p1overlapinterleaved} introduces the generic code of a
+MPI+OpenMP+CUDA implementation, explicitly interleaving
+computation-communication iterations and overlapping MPI communications, CUDA CPU/GPU
+transfers, and CUDA GPU computations.
+
%\begin{algorithm}
% \caption{Generic scheme explicitly overlapping MPI communications, CUDA CPU/GPU transfers and CUDA
% GPU computations, interleaving computation-communication iterations}\label{algo:ch6p1overlapinterleaved}
-\begin{Listing}{algo:ch6p1overlapinterleaved}{Generic scheme explicitly overlapping MPI communications, CUDA CPU/GPU transfers and CUDA GPU computations, interleaving computation-communication iterations}
+\begin{Listing}{algo:ch6p1overlapinterleaved}{generic scheme explicitly overlapping MPI communications, CUDA CPU/GPU transfers, and CUDA GPU computations, interleaving computation-communication iterations}
// Input data and result variables and arrays (example with
// float datatype, 1D input arrays, and scalar results)
float *cpuInputTabAdrCurrent, *cpuInputTabAdrFuture;
cpuResTabAdr = malloc(sizeof(float)*NbIter);
cudaMalloc(&gpuResAdr,sizeof(float));
-// Definition of the Grid of blocks of GPU threads
+// Definition of the grid of blocks of GPU threads
dim3 Dg, Db; Dg.x = ...
// Indexes of source and destination MPI processes
int dest, src; dest = ...
float *gpuFuture = gpuInputTabAdrFuture;
float *tmp;
- // Computation loop on: NbIter + 1 iteration
+ // Computation loop on NbIter + 1 iterations
for (int i = 0; i < NbIter + 1; i++) {
// - Thread 0: achieves MPI communications
if (omp_get_thread_num() == 0) {
sizeof(float)*N, // CPU --> GPU (sync. op)
cudaMemcpyHostToDevice);
}
- // - Thread 2: achieves the GPU computations and the result transfer
+ // - Thread 2: achieves the GPU computations and result transfer
} else if (omp_get_thread_num() == 2) {
if (i > 0) {
gpuKernel_k1<<<Dg,Db>>>(gpuCurrent);// GPU comp. (async. op)
cudaMemcpyDeviceToHost);
}
}
- // - Wait for both threads have achieved their iteration tasks
+ // - Wait until both threads have achieved their iteration tasks
#pragma omp barrier
- // - Each thread permute its local buffer pointers
+ // - Each thread permutes its local buffer pointers
tmp = cpuCurrent; cpuCurrent = cpuFuture; cpuFuture = tmp;
tmp = gpuCurrent; gpuCurrent = gpuFuture; gpuFuture = tmp;
} // End of computation loop
\end{Listing}
%\end{algorithm}
-\Lst{algo:ch6p1overlapinterleaved} introduces the generic code of a
-MPI+OpenMP+CUDA implementation, explicitly interleaving
-computation-communication iterations and overlapping MPI communications, CUDA CPU/GPU
-transfers and CUDA GPU computations. As in the previous algorithms, we declare two CPU input data arrays
-(current and future version) on line 3, but we also declare two GPU input data arrays on line 4. On
-lines 8--11, these four data arrays are allocated, using \texttt{malloc} and
+As in the previous algorithms, we declare two CPU input data arrays
+(current and future version) on line~3. However, in this version we also declare two GPU input data arrays on line~4. On
+lines~8--11, these four data arrays are allocated, using \texttt{malloc} and
\texttt{cudaMalloc}.
-We do not need to allocate page-locked memory space. On lines 23--65 we
-create an OpenMP parallel region, configured to run three threads (see line 21). Lines 26--30 are
+We do not need to allocate page-locked memory space. On lines~23--65 we
+create an OpenMP parallel region, configured to run three threads (see line~21). Lines~26--30 are
declarations of thread local pointers on data arrays and variables (each thread will use its own
-pointers). On line 33, the three threads enter a computation loop of \texttt{NbIter + 1}
+pointers). On line~33, the three threads enter a computation loop of \texttt{NbIter + 1}
iterations. We need to run one more iteration than with previous algorithms.
-Lines 34--41 are the MPI communications, achieved by the thread number $0$. They send the
+Lines~35--41 are the MPI communications, achieved by the thread number $0$. They send the
current CPU input data array to another CPU, and receive the future CPU input data array from
another CPU, like in previous algorithms. But this thread achieves communications only during the
-\emph{first} \texttt{NbIter} iterations. Lines 43--48 are the CPU to GPU input data
-transfers, achieved by thread number $1$. These data transfers are run in parallel of MPI
-communications. They are run during the \emph{first} \texttt{NbIter} iterations, and transfer
-current CPU input data array into the future GPU data array. Lines 50--57
+\emph{first} \texttt{NbIter} iterations. Lines~43--48 are the CPU to GPU input data
+transfers, achieved by thread number $1$. These data transfers are run in parallel with MPI
+communications. They are run during the \emph{first} \texttt{NbIter} iterations and transfer
+current CPU input data array into the future GPU data array. Lines~50--57
correspond to the code run by
-thread number $3$. They start GPU computations, to process the current GPU input data array, and if
+thread number $2$. They start GPU computations, process the current GPU input data array, and if
necessary
-transfer a GPU result at an index of the CPU result array. These GPU computations and result
+transfer a GPU result to an index of the CPU result array. These GPU computations and result
transfers are run during the \emph{last} \texttt{NbIter} iterations: the GPU computations
-have to wait the first data transfer is ended before to start to process any data, and can not run
-during the first iteration. So, the activity of the third thread is shifted of one iteration
-compared to the activities of other threads. Moreover, the address of the current GPU input data
-array has to be passed as a parameter of the kernel call on line 52, in order the GPU threads access
-the right data array. Like in previous algorithms the GPU result is copied at one index of the CPU
-result array, in lines 53--56, but due to the shift of the third thread activity this index is
+have to wait until the first data transfer is ended before starting to process any data and cannot run
+during the first iteration. So, the activity of the third thread is shifted by one iteration
+compared to the activities of the other threads. Moreover, the address of the current GPU input data
+array has to be passed as a parameter of the kernel call on line~52, in order
+for the GPU threads to access
+the right data array. As in previous algorithms the GPU result is copied to one index of the CPU
+result array, in lines~54--56, but due to the shift of the third thread activity this index is
now \texttt{(i - 1)}.
-Line 60 is a synchronization barrier\index{OpenMP!barrier} of the three OpenMP threads, followed by a pointer permutation
-of local pointers on current and future data arrays, on line 62 and 63. Each
+Line~60 is a synchronization barrier\index{OpenMP!barrier} of the three OpenMP threads, followed by a pointer permutation
+of local pointers on current and future data arrays, on line~62 and 63. Each
thread waits for the completion of other
-threads to use the data arrays, and then permutes its data array pointers before to
-enter a new loop iteration.
+threads to use the data arrays, and then permutes its data array pointers before
+entering a new loop iteration.
-This complete overlap of MPI communications and CPU/GPU data transfers and GPU computations, is not
+This complete overlap of MPI communications, CPU/GPU data transfers, and GPU computations is not
too complex to implement, and can be a solution when GPU computations are not adapted to use CUDA
-streams: when GPU computations can not be split in subparts working on independent subsets of input
-data. However, it requires to run one more iterations (a total of \texttt{NbIter + 1}
-iterations). Then, if the number of iterations is very small, it could be more interesting not to
+streams: when GPU computations cannot be split into subparts working on independent subsets of input
+data. However, this requires running one more iteration (a total of \texttt{NbIter + 1}
+iterations). If the number of iterations is very small, it could be more interesting not to
attempt to overlap CPU/GPU data transfers and GPU computations, and to implement \Lst{algo:ch6p1overlapseqsequence}.
\begin{itemize}
\item The first consists of 17 nodes with an Intel Nehalem quad-core processor
- at 2.67Ghz, 6 Gb RAM and an NVIDIA GeForce GTX480 GPU, each.
+ at 2.67Ghz, 6 Gb RAM, and an NVIDIA GeForce GTX480 GPU, each.
\item The second consists of 16 nodes with an Intel core2 dual-core processor at
- 2.67Ghz, 4 Gb RAM and an NVIDIA GeForce GTX580 GPU, each
+ 2.67Ghz, 4 Gb RAM, and an NVIDIA GeForce GTX580 GPU, each.
\end{itemize}
%
-Both clusters have a Gigabit Ethernet interconnection network that is connected
-through a DELL Power Object 5324 switch. The two switches are linked twice,
-insuring the interconnection of the two clusters. The software environment
+Both clusters have a gigabit Ethernet interconnection network that is connected
+through a Dell Power Object 5324 switch. The two switches are linked twice,
+ensuring the interconnection of the two clusters. The software environment
consists of a Linux Fedora 64bit OS (kernel v. 2.6.35), GNU C and C++ compilers
-(v. 4.5.1) and the CUDA library (v. 4.2).
-
+(v. 4.5.1), and the CUDA library (v. 4.2).
%\subsubsection{Validation of the synchronous approach}
\centering
\includegraphics{Chapters/chapter6/curves/gpuSyncOverlap.pdf}
\caption{Experimental performances of different synchronous algorithms computing a
- dense matrix product}
+ dense matrix product.}
\label{fig:ch6p1syncexpematrixprod}
\end{figure}
\label{ch6:p1block-cyclic}
-We have experimented our approach of synchronous parallel algorithms with a classic
+We have tested our approach of synchronous parallel algorithms with a classic
block cyclic algorithm for dense matrix multiplication\index{matrix
- multiplication!block cyclic}. This problem requires to split two input matrices ($A$ and $B$) on a ring of
-computing nodes, and to establish a circulation of the slices of $A$ matrix on the ring ($B$ matrix
+ multiplication!block cyclic}. This problem requires splitting two input matrices ($A$ and $B$) on a ring of
+computing nodes and establishing a circulation of the slices of $A$ matrix on the ring ($B$ matrix
partition does not evolve during all the run). Compared to our generic algorithms, there is no
partial result to transfer from GPU to CPU at the end of each computing iteration. The part of the
-result matrix computed on each GPU is transferred on the CPU at the end of the computation loop.
+result matrix computed on each GPU is transferred onto the CPU at the end of the computation loop.
We have first implemented a synchronous version without any overlap of MPI communications, CPU/GPU
data transfers, and GPU computations. We have added some synchronizations in the native overlapping
two nodes (starting to use MPI communications). But beyond two nodes we get a classical performance
curve.
-Then, we implemented and experimented \Lst{algo:ch6p1overlapnative}, see
+Then, we implemented and tested \Lst{algo:ch6p1overlapnative}, labeled
\emph{ovlp-native} in \Fig{fig:ch6p1syncexpematrixprod}. The native
-overlap of MPI communications with asynchronous run of CUDA kernels appears efficient. When the
+overlap of MPI communications with the asynchronous run of CUDA kernels appears efficient. When the
number of nodes increases the ratio of the MPI communications increases a lot (because the
computation times decrease a lot). So, there is not a lot of GPU computation
time that remains to be
Finally, we implemented \Lst{algo:ch6p1overlapseqsequence}, overlapping MPI communications
with a GPU sequence including both CPU/GPU data transfers and GPU computations,
-see \emph{ovlp-GPUsequence} in \Fig{fig:ch6p1syncexpematrixprod}. From four
-up to sixteen nodes it achieves better performances than \emph{ovlp-native}: we better overlap
-MPI communications. However, this parallelization mechanism has more overhead: OpenMP threads
-have to be created and synchronized. Only for two nodes it is less efficient than the native
+labeled \emph{ovlp-GPUsequence} in \Fig{fig:ch6p1syncexpematrixprod}. From four
+up to sixteen nodes it achieves better performances than \emph{ovlp-native}: the
+overlapping of MPI communications is wider and thus more efficient. However, this parallelization mechanism has more overhead: OpenMP threads
+have to be created and synchronized. With only two nodes it is less efficient than the native
overlapping algorithm. Beyond two nodes, the CPU multithreading overhead seems compensated.
% No, it doesn't need the more implementation of time, but more implementation
% of code :)
chapter = {Optimization methodology for Parallel Programming of Homogeneous or Hybrid Clusters},
publisher = {Saxe-Coburg Publications},
year = {2013},
- OPTkey = {},
- OPTvolume = {},
- OPTnumber = {},
- OPTseries = {},
- OPTtype = {},
- OPTaddress = {},
- OPTedition = {},
- month = feb,
- OPTpages = {},
- note = {to appear},
- OPTannote = {}
+ isbn = {978-1-874672-57-9},
+ url = {http://www.saxe-coburg.co.uk/pubs/future.htm},
+ month = feb
}
@Book{BCC07,
author = {Bahi, J. M. and Contassot-Vivier, S. and Couturier, R.},
- title = {Parallel Iterative Algorithms: from sequential to grid computing},
+ title = {Parallel Iterative Algorithms: From Sequential to Grid Computing},
publisher = {Chapman \& Hall/CRC},
year = {2007},
series = {Numerical Analysis \& Scientific Computing Series},
Computing Systems and Applications (HPCS'2002)},
pages = "90--97",
year = {2002},
- address = {Moncton, Canada},
+ address = {Moncton, New-Brunswick, Canada},
month = jun,
}
booktitle = {8th International Meeting on High Performance Computing for Computational Science, VECPAR'08},
pages = {251--264},
year = {2008},
- address = {Toulouse},
+ address = {Toulouse, France},
month = jun,
}
@InProceedings{Para10,
author = {Contassot-Vivier, S. and Jost, T. and Vialle, S.},
title = {Impact of asynchronism on {GPU} accelerated parallel iterative computations},
- booktitle = {PARA 2010 conference: State of the Art in Scientific and Parallel Computing},
- OPTpages = {--},
+ booktitle = {PARA 2010 Conference: State of the Art in Scientific and Parallel Computing},
+ pages = {43--53},
year = {2010},
address = {Reykjavík, Iceland},
month = jun,
howpublished = {\url{http://www.openmp.org}}
}
-@article{Hoefler08a,
+@inproceedings{Hoefler08a,
author = {T. Hoefler and A. Lumsdaine},
title = {Overlapping Communication and Computation with High Level Communication Routines},
-journal ={Cluster Computing and the Grid, IEEE International Symposium on},
-OPTvolume = {0},
-isbn = {978-0-7695-3156-4},
+booktitle = {IEEE International Symposium on Cluster Computing and the Grid},
year = {2008},
pages = {572-577},
note = "\url{http://doi.ieeecomputersociety.org/10.1109/CCGRID.2008.15}",
publisher = {IEEE Computer Society},
-address = {Los Alamitos, CA, USA},
+address = {Lyon, France},
+isbn = {978-0-7695-3156-4},
}
@Article{Valiant:BSP,
@inproceedings{gustedt:hal-00639289,
AUTHOR = {Gustedt, J. and Jeanvoine, E.},
- TITLE = {{Relaxed Synchronization with Ordered Read-Write Locks}},
+ TITLE = {{Relaxed synchronization with ordered read-write locks}},
BOOKTITLE = {{Euro-Par 2011: Parallel Processing Workshops}},
YEAR = {2012},
MONTH = May,
@article{clauss:2010:inria-00330024:1,
AUTHOR = {Clauss, P.-N. and Gustedt, J.},
- TITLE = {{Iterative Computations with Ordered Read-Write Locks}},
+ TITLE = {{Iterative computations with ordered read-write locks}},
JOURNAL = {{Journal of Parallel and Distributed Computing}},
PUBLISHER = {Elsevier},
VOLUME = {70},
@Book{C11,
- editor = {JTC1/SC22/WG14},
- title = {Programming languages - C},
- publisher = {ISO},
+ editor = {International standardization working group for the programming language C},
+ title = {Programming Languages -- C},
+ publisher = {ISO/IEC},
year = 2011,
- number = {ISO/IEC 9899},
- edition = {Cor. 1:2012}}
+ number = {9899},
+ edition = {Cor. 1:2012}
+}
%% prebreak = \raisebox{0ex}[0ex][0ex]{\ensuremath{\hookleftarrow}},%
%% commentstyle=\textit, numbersep=1em, numberstyle=\tiny, numbers=left,%
%% numberblanklines=false, mathescape, escapechar=@,
- label=#1, caption={#2}}
+ escapechar=@, label=#1, caption={#2}}
}{}
%\def\N{$\mathbb N$ }
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
\chapterauthor{Sylvain Contassot-Vivier}{Université Lorraine, Loria UMR 7503 \& AlGorille INRIA Project Team, Nancy, France.}
\chapterauthor{Stephane Vialle}{SUPELEC, UMI GT-CNRS 2958 \& AlGorille INRIA Project Team, Metz, France.}
-\chapterauthor{Jens Gustedt}{INRIA Nancy -- Grand Est, AlGorille INRIA Project Team, Strasbourg, France.}
+\chapterauthor{Jens Gustedt}{INRIA Nancy--Grand Est, AlGorille INRIA Project Team, Strasbourg, France.}
\chapter{Development methodologies for GPU and cluster of GPUs}
% Glossaire
\section{Glossary}
\begin{Glossary}
-\item[AIAC] Asynchronous Iterations - Asynchronous Communications.
-\item[Asynchronous iterations] Iterative process where each element is updated
+\item[AIAC] Asynchronous Iterations and Asynchronous Communications.
+\item[Asynchronous iterations] iterative process where each element is updated
without waiting for the last updates of the other elements.
-\item[Auxiliary computations] Optional computations performed in parallel to the
+\item[Auxiliary computations] optional computations performed in parallel to the
main computations and used to complete them or speed them up.
-\item[BSP parallel scheme] Bulk Synchronous Parallel, a parallel model that uses
- a repeated pattern (superstep) composed of: computation, communication, barrier.
-\item[GPU stream] Serialized data transfers and computations performed on a same
+\item[BSP parallel scheme] bulk Synchronous Parallel, a parallel model that uses
+ a repeated pattern (superstep) composed of computation, communication, barrier.
+\item[GPU stream] serialized data transfers and computations performed on a same
piece of data.
-\item[Message loss/miss] Can be said about a message that is either not
+\item[Message loss/miss] can be said about a message that is either not
sent or sent but not received (possible with unreliable communication protocols).
-\item[Message stamping] Inclusion of a specific value in messages of the same tag to
+\item[Message stamping] inclusion of a specific value in messages of the same tag to
distinguish them (kind of secondary tag).
-\item[ORWL] Ordered Read Write Locks, a programming tool proposing a unified
+\item[ORWL] Ordered Read-Write Locks, a programming tool proposing a unified
programming model.
-\item[Page-locked data] Data that are locked in cache memory to ensure fast accesses.
-\item[Residual] Difference between results of consecutive iterations in an
+\item[Page-locked data] data that are locked in cache memory to ensure fast accesses.
+\item[Residual] difference between results of consecutive iterations in an
iterative process.
\item[Streamed GPU sequence] GPU transfers and computations performed
simultaneously via distinct GPU streams.
}
@article{ch7:Cavaleri2007603,
-title = "Wave modelling - The state of the art",
+title = "Wave modelling--The state of the art",
journal = "Progress in Oceanography",
volume = "75",
number = "4",
}
@article{ch7:AbottEtAl1978,
-author = "Abott, M. B. Petersens, H. M. and Skovgaard, O.",
+author = "Abott, M. B. and Petersens, H. M. and Skovgaard, O.",
title = "On the numerical modelling of short waves in shallow water",
journal = "Journal of Hydraulic Research",
volume = "16",
@ARTICLE{ch7:MS98,
AUTHOR = "Madsen, P. A. and Sch{\"{a}}ffer, H. A.",
-TITLE = "Higher order Boussinesq-type equations for surface gravity waves - derivation and analysis.",
+TITLE = "Higher order Boussinesq-type equations for surface gravity waves--derivation and analysis.",
JOURNAL = "Advances in Coastal and Ocean Engineering",
VOLUME = "356",
YEAR = "1998",
