\chapter{Cerium TaskManager の GPGPU への対応}
本章では、まずはじめに GPU Programming の問題点について述べ、Cerium TaskManager への実装でそれをどう改善したのかについて説明する。

\section{GPU Programming の問題点}
まず Many Core に対するプログラミングと同様にプログラムの性能を向上させるためにプログラム全体を対象にした並列度を高くしなければならない。
明示的な並列化部分はループ部分である。
GPU は数百個のコアを有しており、ループ部分に対してはデータ並列で処理を行うことで CPU 単体で全てを計算させる場合より圧倒的に高速で演算することができる。
プログラムの大部分がループであれば、データ並列による実行を行うだけでプログラムの性能は十分に向上することになる。しかし、多くのプログラムはその限りではない。
GPU 上での実行においてネックになる部分はデータ転送である。項:\ref{sec:memory_access}で述べたとおり、Shared Memory ではないため host と device でデータの共有ができない。
データを参照するためには Memory 空間のコピーを行うしかない。
これが大きなオーバーヘッドになるので、パイプラインで実行することでデータ転送をオーバーラップする必要がある。
しかし、パイプライン実行の記述は非常に煩雑で、また、パイプラインで実行されていることを証明することは難しくプログラマに大きな負担をかけることになる。
今回、これら問題を解決するため、Cerium TaskManager から受け取った TaskList に従って自動でパイプライン実行するように OpenCL および CUDA で Scheduler を実装した。

\section{OpenCL 実装}
Scheduler と CpuThreads に対応させる形で、GpuScheduler, GpuThreads を実装した。
TaskList から memory 領域を確保し、input data の数だけ clEnqueueWriteBuffer、実行する形式(タスク並列、データ並列)に合わせて clEnqueueTask または clEnqueueNDRangeKernel、output data の数だけ clEnqueueReadBuffer の順序で Command Queue に Operation を発行する。Operation を投入する Command Queue は任意の数に変更することができる。デフォルトでは8段のパイプラインとして実行を行う。
Operation の終了は、clWaitForEvent によって検出し、TaskManger 間の通信を担当する同期キューである mail を使って通知する(図:\ref{fig:createTask})。

GpuScheduler 内で platform\_id や device\_id の取得、Context, Command Queue の生成、device 上での memory 領域の確保、kernel の build と load 等を行なっている。

現在 kernel の記述は、CPU 上で実行する場合と GPU 上で実行する場合はほとんど同じであるが、修飾子など若干形式が異なる。これらは将来的に自動変換など行うのが望ましい。

\section{CUDA 実装}
CUDA の場合も Scheduler と CpuThreads に対応させる形で、CudaScheduler, CudaThreads を実装した。
OpenCL での実装と同様に、TaskList から memory 領域を確保し、input data の数だけ cuMemcpyHtoDAsync、実行する形式(タスク並列、データ並列)に合わせて cuLaunckKernel のブロック数およびスレッド数、次元数を変更し、output data の数だけ cuMemcpyDtoHAsync の順序で Stream に Opration を発行する。Operation を投入する Stream は任意の数に変更することができ、OpenCL での実装と同様に、デフォルトでは8段のパイプラインとして実行を行う。
Operation の終了は、cuStreamQuery によって検出し、終了を待つ必要があるときは cuStreamSynchronize で Stream に発行された全ての Operation の完了を待つ。
TaskManager 間の通信は OpenCL での実装と同様に、同期キューである mail を使って通知する(図:\ref{fig:createTask})。

CUDA 実装も同様に CudaScheduler 内で Context, Stream の生成、device 上での memory 領域の確保、kernel の load 等を行なっている。
Cerium を用いて GPU Programming を行う場合、OpenCL, CUDA のどちらも host 側のコードは同等の記述ができ、プログラマは並列計算のみに集中できる。

kernel の記述は、OpenCL と同様にほとんど同じであるが、修飾子など若干形式が異なる。こちらも CPU 版の kernel から OpenCL 版の kernel などに自動で変換されるようにするのが望ましい。