--- a/paper/Benchmark.tex	Thu Feb 20 02:58:55 2014 +0900
+++ b/paper/Benchmark.tex	Tue Feb 25 19:42:14 2014 +0900
@@ -1,5 +1,32 @@
 \chapter{Benchmark}
+本章では、WordCount, QuickSort を例題として使用し、本研究で実装した GpuScheduler および CudaScheduler の測定を行う。
 
+実験環境
+\begin{itemize}
+\item OS : MacOS 10.9.1
+\item CPU : 2*2.66GHz 6-Core Intel Xeon
+\item GPU : NVIDIA Quadro K5000 4096MB
+\item Memory : 16GB 1333MHz DDR3
+\item Compiler : Apple LLVM version 5.0 (clang-500.2.79) (based on LLVM 3.3svn)
+\end{itemize}
 \section{WordCount}
+WordCount の Task の構成は以下のようになる。
+\begin{itemize}
+\item WordCountTask
+\item DataParallelWordCountTask
+\item PrintTask
+\end{itemize}
+WordCountTask は、input された data を word count し、単語数と行数を output に指定された data 領域に書き込む Task である。
+word count 対象として入力されたファイルは、mmap を用いて memory に展開され、その後データを 16kbyte の大きさに分割して、Task に割り当てられる。
+分割されたデータが送られてくるため、分割された前後のテキストがどうなっているかはわからない。そのため担当範囲であるデータの先頭と末尾のパラメータを単語数と行数の他に付け加える。
+後にそのデータを他の word count 結果と照らし合わせ、分割されたテキストを正しく整合する。
 
-\section{FFT}
+DataParallelWordCountTask も WordCountTask と同様に input された data を word count し、単語数と行数を output に指定された data 領域に書き込む Task である。
+WordCountTask と異なる点は、送られてくるデータが分割されてなく、Task 内で index から担当する範囲を算出し、output に書き込む点である。
+この Task はデータ並列で実行される。GPU 上で実行する場合のデータ転送の回数を減らすことでオーバーヘッドを削減している。
+
+PrintTask は WordCountTask または DataPrallelWordCountTask によって書き出された単語数と行数を集計し、出力する Task である。
+WordCountTask または DataParallelWordCountTask の終了を待ち、動作する。
+
+今回は 
+\section{QuickSort}

--- a/paper/Cerium.tex	Thu Feb 20 02:58:55 2014 +0900
+++ b/paper/Cerium.tex	Tue Feb 25 19:42:14 2014 +0900
@@ -39,6 +39,60 @@
 例えば、Task B には Task A との依存関係がある場合、Task B は Wait Queue に投入され依存関係が解消されるのを待つ。
 Task A は Active Queue に投入され、cpu type によって示された Device に転送され実行される。
 Task A の処理が終了すると Mail で Task B へ通知される。Task B はその通知を受けると Active Queue に投入され、Task A と同様に指定された Device 上で実行される。
+図:\ref{fig:createTask}は Cerium が Task を生成/実行する場合のクラスの構成である。
+\begin{figure}[ht]
+  \begin{center}
+    \includegraphics[scale=0.44]{./images/createTask.pdf}
+  \end{center}
+  \caption{Task Manager}
+  \label{fig:createTask}
+\end{figure}
+
+以下に Task を生成する例題を示す。input data を二つ用意し、input 同士を乗算し、output に格納する multiply という例題である。
+\begin{verbatim}
+void
+multiply_init(HTaskPtr twice, int *i_data, int *o_data) {
+    multiply = manager->create_task(MULTIPLY_TASK);
+    // MULTIPLY_TASK is task id(enum)
+    multiply->set_inData(0, i_data1,
+                         sizeof(int)*length);
+    multiply->set_inData(1, i_data2,
+                         sizeof(int)*length);
+    multiply->set_outData(0, o_data,
+                          sizeof(int)*length);
+    multiply->set_param(0, (memaddr)length);
+    multiply->set_cpu(SPE_ANY);
+    multiply->spawn();
+}
+\end{verbatim}
+
+Task の記述は以下のようになる。表:\ref{table:taskAPI}は Task 側で使用する API である。
+\begin{verbatim}
+static int
+run(SchedTask *s,void *rbuf, void *wbuf)
+{
+    float i_data1=(float*)s->get_input(rbuf,0);
+    float i_data2=(float*)s->get_input(rbuf,1);
+    float o_data=(float*)s->get_output(wbuf,0);
+    long length=(long)s->get_param(0); 
+    for (int i=0;i<length;i++)
+        outdata[i]=indata1[i]*indata2[i];
+    return 0;
+}
+\end{verbatim}
+
+\begin{table}[ht]
+  \begin{center}
+    \label{table:taskAPI}
+    \small
+    \begin{tabular}[t]{|c|l|} \hline
+      get\_input & Schedulerからinput dataを取得 \\ \hline
+      get\_output & Schedulerからoutput dataを取得 \\ \hline
+      get\_param & set\_paramした値を取得 \\ \hline
+    \end{tabular}
+    \caption{ Task 側で使用する API }
+  \end{center}
+\end{table}
 
 \section{Task のスケジューリング}
 Device は、Task を一つずつ受け取るのではなく、ある程度まとめて受け取る。それを TaskList と呼んでいる。TaskList は各 Device 用に TaskManager 側で生成される。
@@ -47,13 +101,14 @@
 
 \begin{figure}[htpb]
   \begin{center}
-    \includegraphics[scale=0.6]{./images/scheduler.pdf}
+    \includegraphics[scale=0.4]{./images/scheduler.pdf}
   \end{center}
   \caption{Task Scheduler}
   \label{fig:task_scheduler}
 \end{figure}
 
 \newpage
+
 \section{Task の入出力}
 Task の入出力の API として、set\_inData, set\_outData がある。
 \begin{description}
@@ -67,3 +122,48 @@
 \item[set\_outData(index, addr, size)] は、Task のデータの出力先を指定する。使用方法は set\_inData と同じで、alignment, byte 数に気をつける必要がある。
   
 \end{description}
+
+\section{Cerium におけるデータ並列}
+Cerium でデータ並列を利用するために、OpenCL の API に合わせた iterate という API を用意した。
+iterate は length を引数として受け取り、Scheduler で length の値と受け取った引数の個数を次元数として Task 数を計算し、データ並列として実行する。
+
+例として、CPU 数4、一次元で10個のデータに対してデータ並列実行を行なった場合、各 CPU が担当する index は表:\ref{table:dpi}のようになる。
+
+\begin{table}[h]
+  \begin{center}
+    \small
+    \begin{tabular}[t]{c||c|c|c|c} \hline
+      stage&CPU0& CPU1&CPU2&CPU3 \\ \hline
+      1&0&1&2&3 \\ \hline
+      2&4&5&6&7 \\ \hline
+      3&8&9& & \\ \hline
+    \end{tabular}
+  \end{center}
+  \label{table:dpi}
+  \caption{Data 並列実行時の index の割り当て}
+\end{table}
+この例だと各 CPU に対する index の割り当ては、CPU0 は index 0,4,8、CPU1 は index 1,5,9、CPU2 は index 2,6、CPU3 は index 3,7 となっている。
+
+\newpage
+\subsection{データ並列による実行}
+Scheduler で、データ並列による実行を行う Task を検出し、各 CPU で Task を実行する、各 CPU が担当する index は SchedTask に格納してある。
+Task は以下のように記述する。
+
+\begin{verbatim}
+static int // Task
+run(SchedTask *s,void *rbuf, void *wbuf)
+{
+    float *indata1,*indata2,*outdata;
+
+    indata1 = (float*)s->get_input(rbuf, 0);
+    indata2 = (float*)s->get_input(rbuf, 1);
+    outdata = (float*)s->get_output(wbuf, 0);
+
+    uisigned long i = s->x;
+    outdata[i]=indata1[i]*indata2[i];
+    return 0;
+}
+\end{verbatim}
+
+並列プログラミングでは、並列化する Task が全部同一であるということは少なくない。
+itrerate を実装したことで、Task を生成する部分をループで回す必要はなくなり、OpenCL と同様に一つの記述で異なる index を持つ Task を CPU 上で実行することができるようになった。

--- a/paper/GPGPU.tex	Thu Feb 20 02:58:55 2014 +0900
+++ b/paper/GPGPU.tex	Tue Feb 25 19:42:14 2014 +0900
@@ -2,9 +2,32 @@
 本章では、まずはじめに GPU Programming の問題点について述べ、Cerium TaskManager への実装でそれをどう改善したのかについて説明する。
 
 \section{GPU Programming の問題点}
-
-\section{パイプライン実行}
+まず 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 などに自動で変換されるようにするのが望ましい。

--- a/paper/Parallel.tex	Thu Feb 20 02:58:55 2014 +0900
+++ b/paper/Parallel.tex	Tue Feb 25 19:42:14 2014 +0900
@@ -11,7 +11,7 @@
 OpenCL C は演算用プロセッサ(本研究では GPU)上で動作する、C 言語を拡張したプログラミング言語である。
 一方で、OpenCL Runtime API は OpenCL C で記述したプログラムを GPU 上で実行させるため、制御用のプロセッサ(本研究では CPU)上で利用する API である。
 
-オペレーティングシステムなどが処理される、メイン CPU などのことを host、GPGPU を搭載したグラフィックボードなどのことを device と定義している。
+OpenCL ではオペレーティングシステムなどが処理されるメイン CPU などのことを host、GPGPU を搭載したグラフィックボードなどのことを device と定義している。
 OpenCL では device に CPU を割り当てることも可能である。OpenCL Application は host 側のプログラムと device 側のプログラムが一体となって動作する。
 この device 上で動作するプログラムを OpenCL では、特別に kernel と呼ぶ。
 
@@ -19,15 +19,15 @@
 OpenCL では、デバイスの操作に Command Queue を使用する。Command Queue は device に OpenCL の Operation を送るために仕組みである。
 Command Queue は clCreateCommandQueue という OpenCL API に所属するコンテキストと実行対象となる device を指定することで生成される。
 
-Command Queue では kernel の実行、input buffer の読み込み、output buffer への書き込みといった Operation が in order で実行される。
-Command Queue を作成するとき CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティを指定することで Operation を out of order で実行することが可能になる。
+Command Queue では kernel の実行、input buffer の読み込み、output buffer への書き込みといった Operation が投入された順序で実行される。
+Command Queue を作成するとき CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティを指定することで Operation を順序を無視して実行することが可能になる。
 Operation を out of order で実行する場合、データの依存関係を記述する必要がある。
 各 Operation には event\_wait\_list と event を指定することができ、これらを利用してデータの依存関係を記述することができる。
 しかし、この CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティをサポートしている device は少なく、Mac OS X では OS レベルでサポートしていない。
 パイプライン実行を行うためには kernel の実行やデータ転送を out of order で実行する必要がある。
 CL\_QUEUE\_OUT\_OF\_ORDER\_EXEC\_MODE\_ENABLE のプロパティが無効の場合、複数の Command Queue を生成し、Command Queue を複数投入することで Operation を out of order で実行することが可能になる。
 
-\subsection{Memory Access}
+\subsection{Memory Access} \label{sec:memory_access}
 host 側は主にデータを input/output する Memory の確保を行う。
 GPU の Memory 空間(図:\ref{fig:gpuarch})や Cell の Memory 空間(図:\ref{fig:cellarch})は Multi Core CPU(図:\ref{fig:cpuarch})とは異なり、Shared Memory ではないため host と device 間でデータの共有ができない。
 アクセスするには Memory 空間ごとコピーしなければならない。
@@ -65,14 +65,14 @@
   
     
 \subsection{Data Parallel Execution}
-多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
+3D グラッフィクのような多次元のデータ構造がある場合に高い並列度を保つには、それを分割して並列に実行する機能が必要である。
 これを OpenCL ではデータ並列と呼んでいる。
 OpenCL では次元数に対応する index があり、OpenCL は一つの記述から index の異なる複数の kernel を自動生成する。
 その添字を global\_id と呼ぶ。このとき入力されたデータは WorkItem という処理単位に分割される。
 
 OpenCL は WorkItem に対して、それぞれを識別する ID(global\_id) を割り当てる。
 kernel は get\_global\_id という API によって ID を取得し、取得した ID に対応するデータに対して処理を行うことでデータ並列を実現する。
-また、WorkItem は3次元までにデータを渡すことができる。
+また、WorkItem は3次元までデータを渡すことができる。
 
 データ並列による kernel 実行の場合、clEnqueueNDRangeKernel API を使用する。この関数の引数として WorkItem の数と次元数を指定することでデータ並列で実行できる。
 
@@ -119,24 +119,26 @@
 例えば get\_global\_id(1) と呼び出した場合は y 座標の、get\_global\_id(2) と呼び出した場合は z 座標の global\_id を取得する。
 
 \section{CUDA}
-CUDA とは、半導体メーカー NVIDIA 社が提供する GPU コンピューティング向けの総合開発環境でプログラム記述、コンパイラ、デバッガなどから構成される。
+CUDA とは、半導体メーカー NVIDIA 社が提供する GPU コンピューティング向けの総合開発環境でコンパイラ、ライブラリ、デバッガなどから構成される。
 プログラム言語である CUDA C は C 言語ベースに拡張を加えたものである。
 
 CUDA には CUDA Runtime API と CUDA Driver API の2種類がある。
 Driver API は Runtime API と比べてプログラマが管理すべきリソースが多い。しかし、Runtime API より柔軟な処理を行うことができる。
 今回は Driver API を使用して実装した。
 
-CUDA も OpenCL と同様に、GPU 側を device、制御を行う CPU 側を host と定義している。また、device 上で動作するプログラムも OpenCL と同様に kernel と呼ぶ。
+CUDA も OpenCL と同様に、制御を行う CPU 側を host、GPU 側を device と定義している。また、device 上で動作するプログラムも OpenCL と同様に kernel と呼ぶ。
 
-\subsection{Stream}
+\subsection{Stream} \label{sec:stream}
 CUDA には OpenCL の Command Queue と似たような仕組みとして Stream がある。
 Stream は host 側の発行された Operation を一連の動作として device で実行する。Stream 内の Operation は発行された順序で実行されることが保証されている。
-異なる Stream での Operation でデータの依存関係などが解消され実行可能な場合、同時に実行することができ、Interleave させることができる。
+異なる Stream での Operation の依存関係が解消され実行可能な場合、Operation を同時に実行することができる。
+例として、ある Stream に kernel を実行する Operation があり、それとは異なる Stream に依存関係がないデータを転送する Operation があった場合、kernel の実行中にデータ転送を行うことが可能になる。
+
 Stream は cuStreamCreate という Driver API で生成される。
 OpenCL と異なり、コンテキストと実行対象となる device を指定する必要はないが、コンテキストを作成した Thread と同一の Thread でないと Stream が生成できないという制約がある。
 引数に Stream を指定しない API はすべて host 側をブロックする同期的な処理となる。複数の Stream を同時に走らせ Operation を並列に実行するためには非同期処理を行う API を利用する必要がある。
 
-Stream 内の Operation を同期を行う方法はいくつかある。
+Stream 内の Operation を同期する方法はいくつかある。
 一つ目は cuStreamSynchronize API を利用した同期方法である。
 cuStraemSynchronize API の引数に Stream を指定すると、指定した Stream に発行されたすべての Operation が終了するまで host をブロックする。
 host をブロックすることなく、Stream に発行された Operation が終了したかどうかを調べるには cuStreamQuery API を利用する。
@@ -147,7 +149,7 @@
 この event は別の Stream で cuEventRecord されるものでも待つことができる。
 
 \subsection{Memory Access}
-CUDA も OpenCL と同様に Shared Memory ではないため host と device 間でデータの共有ができない。アクセスするには Memory 空間ごとコピーする必要がる。
+CUDA も OpenCL と同様に Shared Memory ではないため host と device 間でデータの共有ができない。アクセスするには Memory 空間ごとコピーする必要がある。
 CUDA でのデータの読み込みは cuMemcpyHtoD、書き込みは cuMemcpyDtoH という API でそれぞれ行われる。しかし、これらの API は同期的に実行されてしまう。
 非同期処理にしたい場合、読み込みで cuMemcpyHtoDAsync、書き込みで cuMemcpyDtoHAsync という API をそれぞれ利用することで非同期に行うことができる。
 
@@ -167,23 +169,489 @@
 各組み込み変数はベクター型で、blockDim.x とすると x 座標の Thread 数を参照することができる。
 blockIdx.x とすると x 座標の block ID が参照でき、threadIdx.x とすると x 座標の thread ID を参照することができる。
 blockDim.x * blockIdx.x + threadIdx.x を計算すると OpenCL の get\_global\_id(0) で取得できる ID に相当する ID を得ることができる。
-例えば、ある kernel で get\_global\_id(0) の返り値が13の場合、 CUDA では図:\ref{fig:culculate_index}のようにすることで ID を算出することができる。
+例えば、ある kernel で get\_global\_id(0) の返り値が13の場合、 CUDA では図:\ref{fig:calculate_index}のようにすることで ID を算出することができる。
 
 \begin{figure}[!h]
   \begin{center}
     \includegraphics[scale=0.4]{./images/culculate_index.pdf}
   \end{center}
-  \caption{Culculate Index}
-  \label{fig:culculate_index}
+  \caption{Calculate Index}
+  \label{fig:calculate_index}
 \end{figure}
 
 \newpage
 
 \section{Porting to OpenCL to CUDA}
 本項では OpenCL で記述された Application を CUDA に移植する方法について説明する。
+以下の表は OpenCL と CUDA の用語および修飾子、ID の参照、Object、API の対応表である。
+\begin{table}[!h]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|l|l|} \hline
+      OpenCL & CUDA \\ \hline \hline
+      WorkItem  & Thread  \\ \hline
+      WorkGroup  & Block \\ \hline
+      Global Memory & Global Memory \\ \hline
+      Local Memory & Shared Memory \\ \hline
+      Private Memory & Local Memory \\ \hline
+    \end{tabular}
+    \caption{用語}
+    \label{table:terminology_comp}
+  \end{center}
+\end{table}
 
-\subsection{Sequential Execution}
+\begin{table}[!h]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|l|l|} \hline
+      OpenCL & CUDA \\ \hline \hline
+      \_\_kernel function & \_\_global\_\_ function \\ \hline
+      No necessary & \_\_device\_\_ function (not callable from host) \\ \hline
+      \_\_constant variable & \_\_constant\_\_ variable \\ \hline
+      \_\_global variable & \_\_device\_\_ variable \\ \hline
+      \_\_local variable & \_\_shared\_\_ variable \\ \hline
+    \end{tabular}
+    \caption{修飾子}
+    \label{table:Qualifiers_comp}
+  \end{center}
+\end{table}
+
+\begin{table}[!h]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|l|l|} \hline
+      OpenCL & CUDA \\ \hline \hline
+      get\_num\_groups() & gridDim \\ \hline
+      get\_local\_size() & blockDim \\ \hline
+      get\_group\_id() & blockIdx \\ \hline
+      get\_local\_id() & threadIdx \\ \hline
+      get\_global\_id() & blockDim * blockIdx + threadIdx \\ \hline
+      get\_global\_size() & gridDim * blockDim \\ \hline
+    \end{tabular}
+    \caption{kernel Indexing}
+    \label{table:kernel_api_comp}
+  \end{center}
+\end{table}
+
+\newpage
+
+\begin{table}[!h]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|l|l|} \hline
+      OpenCL & CUDA \\ \hline \hline
+      cl\_device\_id & CUdevice \\ \hline
+      cl\_context & CUcontext \\ \hline
+      cl\_program & CUmodule \\ \hline
+      cl\_kernel & CUfunction \\ \hline
+      cl\_mem & CUdeviceptr \\ \hline
+      cl\_command\_queue & CUstream(but imperfection) \\ \hline
+    \end{tabular}
+    \caption{Objects}
+    \label{table:object_comp}
+  \end{center}
+\end{table}
+
+\begin{table}[!h]
+  \begin{center}
+    \small
+    \begin{tabular}[htpb]{|l|l|} \hline
+      OpenCL & CUDA \\ \hline \hline
+      No required & cuInit() \\ \hline
+      clGetContextInfo() & cuDeviceGet() \\ \hline
+      clCreateContext() & cuCtxCreate() \\ \hline
+      clCreateCommandQueue() & cuStreamCreate() \\ \hline
+      clCreateProgramWithSource() & cuModuleLoad() \\ \hline
+      clBuildProgram() & CUDA programs are compiled offline \\ \hline
+      clCreateKernel() & cuModuleGetFunction() \\ \hline
+      clCreateBuffer() & cuMemAlloc() \\ \hline
+      clEnqueueWriteBuffer() & cuMemcpyHtoD() \\ \hline
+      clEnqueueReadBuffer() & cuMemcpyDtoH() \\ \hline
+      clEnqueueNDRangeKernel() & cuLaunchKernel() \\ \hline
+      clSetKernelArg() & Functonality in cuLaunchKernel() \\ \hline
+      clReleaseMemObj() & cuMemFree() \\ \hline
+    \end{tabular}
+    \caption{APIs}
+    \label{table:api_comp}
+  \end{center}
+\end{table}
+
+\newpage
+
+\subsection{Sequential Execution} \label{sec:seq}
+OpenCL および CUDA で逐次実行するプログラムを例として変換方法を説明する。与えられた二つの input data を乗算し、指定された領域に output する kernel を複数回起動している。
+
+\subsubsection{Initialize}
+\begin{verbatim}
+  // initialize(OpenCL)
+  cl_platform_id platform_id;
+  cl_uint num_platforms;
+  cl_device_id device_id;
+  cl_uint num_devices;
+  cl_uint ret;
+  cl_command_queue command_queue;
+
+  clGetPlatformIDs(1, &platform_id, &num_platforms);
+  clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU,
+                 1, &device_id, &num_devices);
+  cl_context context = clCreateContext(NULL, 1, &device_id,
+                                       NULL, NULL, &ret);
+  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
+\end{verbatim}
+
+OpenCL での初期化は上記のようになる。
+
+OpenCL は様々なメーカー(NVIDIA, AMD など) GPU の対応しているため platform\_id を取得し、それをもとに Context を生成する。
+
+\begin{verbatim}
+  // initialize(CUDA)
+  CUdevice device;
+  CUcontext context;
+  
+  cuInit(0);
+  cuDeviceGet(&device, 0);
+  cuCtxCreate(&context, CU_CTX_SCHED_SPIN, device);
+\end{verbatim}
+
+CUDA での初期化は上記のようになる。
+
+CUDA は NVIDIA が提供する GPU にのみ対応しているので個別の id を取得する必要がない。
+また、CUDA には Null Stream と呼ばれるデフォルトの Stream がある。
+OpenCL の Command Queue のように必ず生成する必要はない。
+\subsubsection{Load Kernel}
+\begin{verbatim}
+  // load kernel(OpenCL)
+  const char* filename = "multiply.cl";
+  const char* functionname = "multiply";
+  
+  int fp = open(filename, O_RDONLY);
+  
+  struct stat stats;
+  fstat(fp,&stats);
+  off_t size = stats.st_size;
+  
+  char *kernel_src_str = (char*)alloca(size+1);
+  size_t kernel_code_size = read(fp, kernel_src_str, size);
+  close(fp);
+  kernel_src_str[size] = 0;
+
+
+  cl_program program;
+  program = clCreateProgramWithSource(context, 1,
+                                      (const char **)&kernel_src_str, 0, &ret);
+  clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
+  cl_kernel kernel = clCreateKernel(program,functionname, &ret);
+\end{verbatim}
+
+OpenCL での kernel の読み込みは上記のようになる。
+
+OpenCL ではプログラム内で kernel をコンパイルすることが可能である。
+専用のコンパイラを必要としないため、様々なプラットフォームへの展開が可能になる。
+また、Linux 環境では Fixstars 社が提供する foxc、Mac OS X では openclc を利用することで事前にコンパイルすることもできる。
+
+\begin{verbatim}
+  // load kernel(CUDA)
+  CUmodule module;
+  CUfunction function;
+
+  cuModuleLoad(&module, "multiply.ptx");
+  cuModuleGetFunction(&function, module, "multiply");
+\end{verbatim}
+
+CUDA での kernel の読み込みは上記のようになる。
+
+CUDA では CUDA に付属されている専用コンパイラ nvcc を使って事前に kernel をコンパイルする必要がある。
+
+\subsubsection{Memory Allocate}
+\begin{verbatim}
+  // memory allcate(OpenCL)
+  cl_mem memA = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                               WORKS*sizeof(float), NULL, &ret);
+  cl_mem memB[num_exec];
+  cl_mem memOut[num_exec];
+  for (int i=0;i<num_exec;i++) {
+    memB[i] = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                             sizeof(float), NULL, &ret);
+    memOut[i] = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
+                               WORKS*sizeof(float), NULL, &ret);
+  }
+\end{verbatim}
+
+OpenCL での device 上の memory 領域の確保は上記のようになる。
+
+今回のプログラムでは一つ目の input data は共通の値なので、一つ目の input data を格納する領域は一つのみ確保している。
+二つ目の input data および output data は kernel ごとに異なる値なので、実行する kernel の数だけ領域を確保している。
+
+\begin{verbatim}
+  // memory allocate(CUDA)
+  CUdeviceptr devA;
+  CUdeviceptr devB[num_exec];
+  CUdeviceptr devOut[num_exec];
+  
+  cuMemAlloc(&devA, LENGTH*THREAD*sizeof(float));
+  for (int i=0;i<num_exec;i++) {
+    cuMemAlloc(&devB[i], sizeof(float));
+    cuMemAlloc(&devOut[i], LENGTH*THREAD*sizeof(float));
+  }
+\end{verbatim}
+
+CUDA での device 上の memory 領域の確保は上記のようになる。
+
+CUDA では memory 領域を確保するとき、context を生成した Thread と同一の Thread である必要があるため引数に context を指定する必要はない。
+\newpage
+\subsubsection{Data Transfer(host to device)}
+\begin{verbatim}
+  // data transfer(OpenCL)
+  clEnqueueWriteBuffer(command_queue, memA, CL_TRUE, 0,
+                       WORKS*sizeof(float), A, 0, NULL, NULL);
+
+  for (int i = 0;i<num_exec;i++){
+    B[i] = (float)(i+1);
+    clEnqueueWriteBuffer(command_queue, memB[i], CL_TRUE, 0,
+    sizeof(float), &B[i], 0, NULL, NULL);
+  }
+\end{verbatim}
+
+OpenCL での host から device へのデータ転送は上記のようになる。
+
+clEnqueueWriteBuffer の第三引数に CL\_TRUE を指定することで、データ転送が完了するまで host をブロックする同期的な Operation となる。
+\begin{verbatim}
+  // data transfer(CUDA)
+  cuMemcpyHtoD(devA, A, LENGTH*THREAD*sizeof(float));
+  
+  for (int i=0;i<num_exec;i++) {
+    B[i] = (float)(i+1);
+    cuMemcpyHtoD(devB[i], &B[i], sizeof(float));
+  }
+\end{verbatim}
+
+CUDA での host から device へのデータ転送は上記のようになる。
 
+cuMemcpyHtoD はデータ転送が完了するまで host をブロックする同期的な Operation である。
+\subsubsection{Launch Kernel}
+\begin{verbatim}
+  // launch kernel(OpenCL)
+  clSetKernelArg(kernel, 0, sizeof(cl_mem), &memA);
+  for (int i = 0;i<num_exec;i++){
+    clSetKernelArg(kernel, 1, sizeof(cl_mem), &memB[i]);
+    clSetKernelArg(kernel, 2, sizeof(cl_mem), &memOut[i]);
+    clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
+                           &num_work, NULL, 0, NULL, NULL);
+    }
+\end{verbatim}
+
+OpenCL での kernel の実行は上記のようになる。
+
+clSetKernelArg で kernel の引数に対して値を渡すことができる。
+
+clEnqueueNDRangeKernel の第三引数に WorkItem の次元数を指定し、第五引数に WorkItem の数を指定することでデータ並列で実行される。
+\begin{verbatim}
+  // launch kernel(CUDA)
+  for (int i=0;i<num_exec;i++) {
+    void* args[] = {&devA, &devB[i], &devOut[i]};
+    cuLaunchKernel(function,
+                   LENGTH, 1, 1,
+                   THREAD, 1, 1,
+                   0, 0, args, NULL);
+  }
+\end{verbatim}
+
+CUDA での kernel の実行は上記のようになる。
+
+clLaunchKernel の第二、第三、第四引数に各次元のブロック数を指定し、第五、第六、第七引数に各次元の Thread 数を指定することでデータ並列で実行される。
+今回は、LENGTH 個のブロックがあり、各ブロックに THREAD 個の Thread が起動する。
+また、CUDA には clSetKernelArg に相当する API はない。clLaunchKernel の第十引数に指定することで kernel の引数に対して値を渡すことができる。
+\subsubsection{Data Transfer(device to host)}
+\begin{verbatim}
+  // data transfer(OpenCL)
+  for (int i=0;i<num_exec;i++)
+    clEnqueueReadBuffer(command_queue, memOut[i], CL_TRUE, 0,
+                        WORKS*sizeof(float), result[i], 0, NULL, NULL);
+\end{verbatim}
+
+OpenCL での device から host へのデータ転送は上記のようになる。
+
+clEnqueueWriteBuffer と同様に第三引数に CL\_TRUE を指定することでデータ転送が完了するまで host をブロックする同期的な Operation となる。
+\begin{verbatim}
+  // data transfer(CUDA)
+  for (int i=0;i<num_exec;i++)
+    cuMemcpyDtoH(result[i], devOut[i], LENGTH*THREAD*sizeof(float));
+\end{verbatim}
+
+CUDA での device から host へのデータ転送は上記のようになる。
+
+cuMemcpyDtoH は cuMemcpyHtoD と同様にデータ転送が完了するまで host をブロックする同期的な Operation である。
+\newpage
+\subsubsection{Memory Release}
+\begin{verbatim}
+  // memory release(OpenCL)
+  clReleaseMemObject(memA);
+  for (int i=0;i<num_exec;i++) {
+    clReleaseMemObject(memB[i]);
+    clReleaseMemObject(memOut[i]);
+  }
+  clReleaseCommandQueue(command_queue);
+  clReleaseKernel(kernel);
+  clReleaseProgram(program);
+  clReleaseContext(context);
+\end{verbatim}
+
+OpenCL での memory 領域の開放やオブジェクトの開放は上記のようになる。
+
+\begin{verbatim}
+  // memory release(CUDA)
+  cuMemFree(devA);
+  for (int i=0;i<num_exec;i++) {
+    cuMemFree(devB[i]);
+    cuMemFree(devOut[i]);
+  }
+  cuModuleUnload(module);
+  cuCtxDestroy(context);
+\end{verbatim}
+
+CUDA での memory 領域の開放やオブジェクトの開放は上記のようになる。
+
+\newpage
 \subsection{Parallel Execution}
+項:\ref{sec:seq}で例として用いたプログラムは一つ目の input data 以外に依存関係が存在しない。
+つまり、OpenCL では Command Queue、CUDA では Stream を複数生成し、Operation を並列に実行して良い。
+Operation を並列に実行することでデータ転送をオーバーラップすることができる。
+このプログラムを kernel を実行する数だけ Command Queue および Stream を生成し、並列実行されるように変更する。
+変更がある部分のみ記載する。それ以外の部分は項:\ref{sec:seq}と同じである。
+実行の様子は図:\ref{fig:parallel_exec}のようになる。
+\begin{figure}[!htpd]
+  \begin{center}
+    \includegraphics[scale=0.5]{./images/paralle_exec.pdf}
+  \end{center}
+  \caption{Parallel Execution}
+  \label{fig:parallel_exec}
+\end{figure}
+
+\subsubsection{Initialize}
+\begin{verbatim}
+  // initialize(OpenCL)
+  cl_command_queue command_queue[num_exec];
+  
+  for (int i=0;i<num_exec;i++)
+    command_queue[i] = clCreateCommandQueue(context, device_id, 0, &ret);
+\end{verbatim}
+
+OpenCL では上記のように変更する。
+
+Command Queue を配列として宣言し、要素の数だけ Command Queue を生成する。
+\begin{verbatim}
+  // initialize(CUDA)
+  CUstream stream[num_exec];
+  
+  for (int i=0;i<num_exec;i++)
+    cuStreamCreate(&stream[i], 0);
+\end{verbatim}
+
+CUDA では上記のように変更する。
+
+Stream を生成するとき、OpenCL と同じように context と device\_id を渡す必要はない。
+これは項:\ref{sec:stream}で述べたとおり、コンテキストを作成した Thread と同一の Thread でないと Stream を生成できないからである。
+
+OpenCL と同様に Stream を配列として宣言し、要素の数だけ Stream を生成する。
+
+\subsubsection{Data Transfer(host to device)}
+\begin{verbatim}
+  // data transfer(OpenCL)
+
+  for (int i = 0;i<num_exec;i++){
+    B[i] = (float)(i+1);
+    clEnqueueWriteBuffer(command_queue[i], memB[i], CL_FALSE, 0,
+    sizeof(float), &B[i], 0, NULL, NULL);
+  }
+\end{verbatim}
+
+OpenCL では上記のように変更する。
+
+clEnqueueWriteBuffer の第三引数に CL\_FALSE を指定することで、非同期な Operation として実行できる。つまり、Command Queue に Operation を投入し、投入した Operation の完了を待たずに host に制御を返す。
+\begin{verbatim}
+  // data transfer(CUDA)
+  
+  for (int i=0;i<num_exec;i++) {
+    B[i] = (float)(i+1);
+    cuMemcpyHtoDAsync(devB[i], &B[i], sizeof(float), stream[i]);
+  }
+\end{verbatim}
+
+CUDA では上記のように変更する。
+
+cuMemcpyHtoDAsync は非同期な Operation として実行される。第四引数に関連する Stream を指定することができる。
+
+\subsubsection{Launch Kernel}
+\begin{verbatim}
+  // launch kernel(OpenCL)
+  for (int i = 0;i<num_exec;i++){
+    clEnqueueNDRangeKernel(command_queue[i], kernel, 1, NULL,
+                           &num_work, NULL, 0, NULL, NULL);
+    }
+\end{verbatim}
+
+OpenCL では上記のように変更する。
+
+関連する Command Queue を指定するだけで良い。
+\begin{verbatim}
+  // launch kernel(CUDA)
+  for (int i=0;i<num_exec;i++) {
+    cuLaunchKernel(function,
+                   LENGTH, 1, 1,
+                   THREAD, 1, 1,
+                   0, stream[i], args, NULL);
+  }
+\end{verbatim}
+
+CUDA では上記のように変更する。
+
+第九引数に関連する Stream を指定するだけで良い。
+\subsubsection{Data Transfer(device to host)}
+\begin{verbatim}
+  // data transfer(OpenCL)
+  for (int i=0;i<num_exec;i++)
+    clEnqueueReadBuffer(command_queue[i], memOut[i], CL_FALSE, 0,
+                        WORKS*sizeof(float), result[i], 0, NULL, NULL);
+\end{verbatim}
+
+OpenCL では上記のように変更する。
+
+clEnqueueWriteBuffer と同様に第三引数に CL\_FALSE を指定することで非同期な Operation となる。
+\begin{verbatim}
+  // data transfer(CUDA)
+  for (int i=0;i<num_exec;i++)
+    cuMemcpyDtoHAsync(result[i], devOut[i], LENGTH*THREAD*sizeof(float), stream[i]);
+\end{verbatim}
+
+CUDA では上記のように変更する。
+
+cuMemcpyDtoHAscyn は cuMemcpyHtoDAsync と同様に非同期な Operation として実行される。
+
+\subsubsection{Wait for Exec}
+host をブロックしない非同期な Operation はどのタイミングで実行されるかわからない。
+全ての Operation が完了していないのに memory やオブジェクトを開放してしまう場合がある。
+その場合、処理が正しく行われない。
+この問題を回避するために、全ての Operation を待つという処理を host 側に記述する必要がある。
+\begin{verbatim}
+  // wait for exec(OpenCL)
+  for (int i=0;i<num_exec;i++)
+    clFinish(command_queue[i]);
+\end{verbatim}
+
+OpenCL での投入された Operation の完了を待つのは上記のようになる。
+
+clFinish に完了を待ちたい Command Queue を指定することで、投入された Operation が全て完了するのを待つことができる。
+\newpage
+\begin{verbatim}
+  // wait for exec(CUDA)
+  for (int i=0;i<num_exec;i++)
+      cuStreamSynchronize(stream[i]);
+\end{verbatim}
+
+CUDA での投入された Operation の完了を待つのは上記のようになる。
+
+OpenCL と同様に cuStreamSynchronize に完了を待ちたい Steram を指定することで、Operation の終了を待つことができる。
+cuStreamSynchronize は指定された Stream が完了するまで host をブロックする。
+完了したかどうかだけを知りたい場合は、cuStreamQuery を利用する。
 
 \section{StarPU}

Binary file paper/thesis-paper.pdf has changed