%----------------------------------------
% -begin runtime-design.tex
%----------------------------------------
%
As noted in the introduction and in \cite{Qingyu-TG11}, Uintah is a sophisticated computational framework that
can integrate multiple simulation components, analyze the dependencies
and communication patterns between them, and execute the
resulting multi-physics simulation. This is done by  utilizing an abstract
task-graph representation of parallel computation and
communication to express data dependencies between components.
The task-graph is a directed acyclic graph of tasks. Each task
consumes some input and produces some output (which is in turn the
input of some future task). These inputs and outputs are specified
for each patch in a structured grid.

Associated with each task is a C++ method which is used to perform the actual computation.
In the context of the new hybrid CPU-GPU scheduler, a GPU task is represented by an additional
C++ method that is used for GPU kernel setup and invocation.
Each component specifies a list of tasks to be performed and the data
dependencies between them. The task-graph approach of Uintah shares many features
with the migratable object philosophy of Charm++ \cite{CharmPP}. In order to increase efficiency,
the task graph is created and stored locally \cite{TGRID10MB}.
Uintah's CPU-GPU task scheduler is responsible for computing the dependencies of tasks,
determining the order of execution and ensuring that the correct inter-process communication
is performed \cite{TGRID10MB}. It also ensures that no input or output variable
conflicts will exist in any two simultaneously running tasks.

In the migration of the Uintah Computational Framework to hybrid CPU-GPU architectures, we elected to 
use Nvidia CUDA C/C++ for numerous reasons. Looking at the upgrade path of the DOE Jaguar XK6 system to
Titan \cite{titandev_webpage} and also the Keeneland Initial Delivery System (KIDS) \cite{keeneland_webpage}, we
see a trend in the use or planned use of Nvidia GPUs. These are the target machines on which we are already
running both CPU and mixed CPU-GPU simulations.
Initial runs using ported portions of the CFD component ICE,
have demonstrated the ability of our CPU-GPU scheduler to run capability jobs on both KIDS and TitanDev,
utilizing all CPU cores and all GPUs simultaneously on each machine. KIDS currently has 1440 CPU cores
and 360 Nvidia Tesla 20-series GPUs and TitanDev, 15360 CPU cores and 960 Nvidia Tesla 20-series GPUs.

The principal additions made by our new CPU-GPU scheduler are:
significant leveraging of the Nvidia CUDA Asynchronous API \cite{nvidiadev-webpage} to
best overlap PCIe transfers and MPI communication with GPU and CPU computation;
insulating the component developer from the complexities and details involved
with device memory management and asynchronous operations, by automatically managing these operations;
using knowledge of the task-graph and task dependencies to pre-fetch data needed for simulation
variables prior to task execution. Hence, when a GPU task is ready to run, data needed for the task
is already resident in GPU main memory. The GPU task need merely query the scheduler
for device pointers and invoke the kernel.

The existing Uintah code base is nearly 700K lines of code, a significant challenge to port in terms of infrastructure
and existing simulation components. Although OpenCL \cite{opencl-webpage} has the potential to support more than just GPUs
and will be a consideration for use in the future, Nvidia CUDA currently offers far greater
support in terms of performance and analysis tools as well as an API allowing for easier performance gains and
portability for existing codes. Below we describe the design of our CPU-GPU scheduler and its use
of the Nvidia CUDA Asynchronous API \cite{nvidiadev-webpage} in detail.

\subsection{Multi-Threaded Runtime System}
The overall design of the multi-threaded MPI runtime system is explained in great detail in \cite{Qingyu-TG11}, but
to provide context, we review its design briefly here. We then describe in detail how this architecture has been extended
by our recent work, adapting Uintah to run on current and emerging heterogeneous systems.

As mentioned in\cite{Qingyu-TG11}, the core scheduler component that stores simulation variables is the \emph{data warehouse}.
The \emph{data warehouse} is a hashed-map-based dictionary which maps a variable name and patch ID to a memory address.
In the Uintah framework, after the regridder changes the simulation grid and the load balancer generates the patch 
distribution,  the scheduler will create new sets of detailed tasks, compile a new task graph and initialize
the \emph{data warehouses}. Uintah's innovative load balancer utilizes space-filling curves in order
to cluster patches together \cite{luitjens1}.
Originally, Uintah used both dynamic and static schedulers, based solely on MPI, in which data structures were created on each MPI process.
Although most of Uintah's infrastructure components were carefully designed to be stored in a distributed manner,
it was necessary for some data to be stored multiple times, e.g. neighboring patch sets, neighboring tasks and ghost variables.
A limitation of pure MPI scheduling was that tasks which were created and executed on the same node could not share data.
Uintah's multi-threaded MPI scheduler \cite{Qingyu-TG11} solves this problem by dynamically assigning tasks to worker threads 
during execution and shares the same infrastructure components between threads. This design uses one control thread and
several worker threads per MPI process. The control thread holds all infrastructure components such as the regridder,
the load balancer, the task graph and the \emph{data warehouse} and has read and write access to them.

Central to the design of the dynamic CPU-GPU scheduler (Figure ~\ref{fig:hybrid-scheduler})
is the multi-stage queuing architecture for efficient scheduling of CPU and GPU tasks.
The CPU-GPU scheduler utilizes four task queues: an internal ready queue and an external ready queue
for CPU tasks and two queues for the GPU; one for initially ready GPU tasks; those that have requisite simulation
variable data copies from host-to-device pending, and a second for the corresponding device-to-host data copies
pending completion. First, if a task's internal
dependencies are satisfied, then that task will be put in the CPU internal ready queue 
where it will wait until all required MPI communication has finished.
In this same step, if the task is GPU-enabled, the task is then put into the host-to-device
copy queue for advancement toward execution. Ultimately, the task goes to the pending device-to-host copies queue.
As long as the CPU external queue is not empty, there are always tasks to run.
Execution of a task takes place on the first available CPU core or GPU and the scheduler
resides on a single, dedicated core per node. CPU tasks are dispatched by the control thread to available CPU
cores when they signal the need for work. GPU tasks are assigned in a round-robin fashion to available GPUs on-node
once their asynchronous host-to-device data copies have completed.
This design helps to overlap MPI communication and asynchronous GPU data transfers with CPU
and GPU task execution, significantly reducing MPI wait times.

\begin{figure*}[htpb]
\centerline{
  \includegraphics[width=0.78\textwidth,draft=false]{gpu_scheduler.png}}
\vspace{-.1in}
  \caption{Uintah CPU-GPU Task Scheduler Architecture}
  \label{fig:hybrid-scheduler}
\vspace{-.15in}
\end{figure*}

\begin{algorithm}
\caption{GPU Task Controller Algorithm}
\label{control}
\begin{algorithmic}
\WHILE{$~ doneTasks ~ < ~ totalTasks$}
\IF {$~ numExternalReadyTasks~() ~ > ~0$}
\IF {$~ highest\ priority\ task\ isGPUEnabled~()$}
  \STATE $initiateH2DCopies~(task, iteration)$
  \STATE $task->markInitiated~()$
  \STATE $addInitiallyReadyGPUTask~(task)$
\ENDIF
\ENDIF
\IF {$~ numInitiallyReadyGPUTasks~() ~ > ~0$}
  \IF {$task->checkH2DCopyDependencies~()$}
    \STATE $runGPUTask~(task, iteration)$
    \STATE $addCompletionPendingGPUTask~(task)$
  \ENDIF
\ENDIF
\IF {$numCompletionPendingGPUTasks() ~ > ~0$}
  \IF {$task->checkD2HCopyDependencies~()$}
    \STATE $postMPISends~(task, iteration)$
    \STATE $reclaimStreams~(task);$
    \STATE $reclaimEvents~(task);$
    \STATE $task->completed~()$
  \ENDIF
\ENDIF
\ENDWHILE
\end{algorithmic}
\end{algorithm}

\subsection{Asynchronous GPU Techniques}
Significant difficulties arise when mixing concurrency APIs, most notably race conditions, deadlock and
general synchronization complexities.
Within Uintah's CPU-GPU scheduler is a combination of MPI, Pthreads and Nvidia CUDA, a combination
that must be managed with care to avoid such difficulties. Multiple GPUs per node further complicate this situation
in the presence of asynchronous memory copies and multiple device contexts (one CUDA calling context per device per process).
In the same fashion that Uintah insulates the application developer from the parallelism its infrastructure
provides, it also hides and carefully manages details related to GPU memory allocation and transfer.
The Fermi-based GPUs found on the target machines mentioned at the beginning of this section offer additional
ways to achieve asynchronous concurrent execution of kernels. These GPUs have two copy engines and
support multiple kernels running concurrently. Using these features, GPU tasks can be copying data
to-and-from the device as well as running multiple kernels simultaneously. In order to exploit these features, the CPU-GPU scheduler
creates and manages queues of CUDA \emph{Streams} \cite{nvidiadev-webpage}, one for each device on-node.
\emph{Streams} provide a means to perform multiple operations simultaneously in that operations from different
\emph{streams} can be interleaved and also run concurrently. Our implementation also uses CUDA \emph{Events} \cite{nvidiadev-webpage}
which are used for timing and in checking completion of operations such as asynchronous memory copies to-and-from the GPU.

\subsection{Extending the Uintah Task Class}
Previously, the portion of the Uintah Task class responsible for actual execution of the C++ method representing
the computation to perform, was comprised of a single instance of an Action class, which contains a single function pointer to the C++ method
to run. With the addition of GPU tasks, we have modified the Uintah Task class to include an additional
Action instance with associated pointer to the function containing the GPU kernel setup and invocation. This has been accomplished
without altering any existing interface or simulation component.

The design decision to support registration of multiple function pointers was to ultimately add the ability for the
scheduler to chose between execution of the CPU or GPU version of the task at runtime. It may be the case
that if all on-node GPUs are currently busy or unavailable and there exists
an idle CPU core, then it is best to execute a particular task on that CPU core. Currently, if a GPU task has a
GPU implementation, it is executed on the GPU. This overall infrastructure design remains broad
enough to use other accelerator designs, such as the Intel MIC \cite{intel-mic-webpage} chip.

\subsection{Pre-fetching GPU Task Data}
When the CPU-GPU scheduler begins dispatching ready tasks from the CPU external ready queue, it diverts GPU-enabled
tasks to the initially-ready GPU task queue. Just prior to this step the CPU-GPU scheduler initiates the
device memory allocations and asynchronous host-to-device data copies for the requisite simulation variables.
This is accomplished by querying the \emph{data warehouse} for the location and size of the data required
for computation and also requesting that the \emph{data warehouse}
allocate space for the result of the computation. We have exposed a flat representation of the underlying 3D data structure
representing each simulation variable on a patch. This linear array maps relatively easily onto the GPU.
To fully exploit the aforementioned levels of concurrency, the 
host memory to be copied to device must be page-locked. This guarantees the memory will not be paged to disk. The CPU-GPU scheduler
then registers for DMA the host memory to be copied to the GPU using a call to \verb+cudaHostRegister()+
with the \verb+cudaHostRegisterPortable+ flag.
This call and flag pair create page-locked (often referred to as pinned) memory from pre-allocated host memory that is
considered page-locked by all CUDA contexts. This step avoids a bounce buffer and accelerates PCIe transfers and also eliminates resetting of CUDA contexts
when referencing the registered host memory. A call to \verb+cudaHostRegister()+ can be cleanly performed from the host without setting
a context.

The new scheduler infrastructure maintains a set of queues for \emph{stream} and \emph{event} handles (one per device representing separate contexts for each),
and assigns them to each simulation variable per time step to overlap with other host-to-device memory copies as well as kernel
execution. These \emph{stream} and \emph{event} handles are stored by the associated Task itself and effectively provide a mechanism to detect
completion of asynchronous memory copies without a busy wait, using \verb+cudaEventQuery(event)+. This allows
querying the status of all device work preceding the most recent CUDA 4.0 API call to \verb+cudaEventRecord()+ \cite{nvidiadev-webpage}.

On systems with multiple on-node GPUs such as KIDS, the CPU-GPU scheduler must also manage a CUDA calling context for each device.
This is set per device prior to subsequent CUDA API calls on that device. In general, the CPU-GPU scheduler assigns a device to the task itself
(round-robin), allocates space on the device, marks the task as initiated and then starts the asynchronous host-to-device memory copies.
The entire GPU task processing algorithm is shown in Algorithm~\ref{control}, where it should be noted that CPU task processing as shown in
\cite{Qingyu-TG11}, is interleaved with the GPU task processing.

A call to \verb+cudaEventRecord()+ is then made after a call to\\* \verb+cudaMemcpyAsync()+ and these \emph{event} pointers are stored
with the task itself, and the task is placed into the initially-ready GPU task queue. Priority of GPU tasks is based
on the same prioritization algorithm used in the CPU external ready queue, thus the overall task priority is preserved.
This is all accomplished asynchronously with respect to the CPU, which is continually responding to requests from
idle CPU cores for work. This series of steps essentially prepares the GPU memory needed by the task and is all
completed prior to task execution. All data related to each task's host and device pointers are kept in a set of maps
maintained by the CPU-GPU scheduler. These maps will ultimately become a separate GPU \emph{data warehouse} in future work.

\subsection{GPU Task Execution}
During successive iterations of the CPU-GPU scheduler's task controller algorithm, the scheduler checks for
existing tasks in the initially-ready GPU task queue and determines if its host-to-device
memory copies have completed. 
This is accomplished by performing \verb+cudaEventQuery(event)+ on each of its recorded \emph{events}.
The scan is essentially linear in the size of the list of \emph{events} to query, but this size is never greater than say 10
elements, and is essentially constant time, $O(1)$. If all \emph{event} queries return with \verb+cudaSuccess+, the GPU task is
ready to run. The C++ method associated with the kernel setup and invocation can then be executed. The component
queries the scheduler for device pointers, and a \emph{stream} to associate with the kernel launch. The
component then passes these to the kernel routine that performs the computation on the device. To transfer the results of the computation
back to the host, the component code requests a device-to-host copy via the infrastructure API.
The scheduler in turn initiates the
asynchronous memory copy from device to host destination and records the \emph{events} associated with the task. Afterward, the task
is placed in the completion pending GPU task queue.

\subsection{GPU Task Completion and MPI Sends}
Within the CPU-GPU scheduler's task processing loop (Algorithm~\ref{control}),
the \emph{events} in the \emph{stream} associated with the device-to-host memory copy (and kernel used to compute results) of the highest
priority GPU task can be queried for completion.
Success returned on each of a task's \emph{events} indicates the task has completed execution. The results are then guaranteed to be in the host-side
\emph{data warehouse}. At this point, the task can be marked as completed and the CPU-GPU scheduler then reclaims all of the
\emph{events} and \emph{streams} used by the task. MPI sends from the GPU task can then be posted. The GPU task is finally removed from the completion pending
task queue, allowing other dependent tasks to proceed.
% 
%----------------------------------------
% -end runtime-design.tex
%----------------------------------------