--- a/paper/chapter5.tex	Wed Feb 04 16:38:20 2015 +0900
+++ b/paper/chapter5.tex	Mon Feb 09 11:32:28 2015 +0900
@@ -3,21 +3,154 @@
 現在、GPU のような異なる種類のアーキテクチャを搭載した CPU 、つまりヘテロジニアスな CPU が増えている。
 特定の計算に特化した Task の生成やスケジューリングを行い、 GPGPU により高い並列度を出す研究は様々な分野で行われている。
 本研究では Cerium を特定の計算に限らず、 GPU を用いて汎用計算を行えるフレームワークに改良する。
-\section{OpenCL による実装}
-\subsection{ベンチマーク}
-OpenCL を用いた GPU 対応を行った。
-Scheduler と CpuThreads に対応させる形で GpuScheduler と GpuThreads を実装した。
-GpuScheduler 内で OpenCL の API を用いて GPU の制御を行っている。
-
-受け取った TaskList からメモリバッファを作成し、 
-clEnqueueuWriteBuffer、clEnqueueTask、clEnqueueReadBuffer の順に CommandQueue に Enqueueしていく。
-Task の投入は CommandQueue を2つ用意しパイプライン的に実行を行う。
-
-Task の終了は、clWaitForEvent によって検出し、TaskManager 間の通信を担当する同期キューである mail を使って通知する。
-
-\section{CUDA による}
-\subsection{ベンチマーク}
 
 
+\section{OpenCL および CUDA による実装}
+OpenCL 、CUDA による GPGPU 対応を行った。
+Scheduler と CpuThreads に対応させる形で
+OpenCL を用いた GpuScheduler と GpuThreads、
+CUDA を用いた CudaScheduler と CudaThreads を実装した。
+それぞれの Scheduler 内で各フレームワークの API を用いて GPU の制御を行っている。
+
+ TaskManager から受け取った TaskList をもとに Device 上のメモリバッファを作成する。
+その後 CommandQueue、 Stream といったそれぞれの Queue に Device 制御用の Command を Queueing していく。
+
+Command は Queueing した順に実行されるので、以下のように Command を Queueing する。
+\begin{enumerate}
+\item Host から Device へのデータ転送
+\item kernel の実行
+\item Device から Host へのデータ転送
+\end{enumerate}
+
+データの転送や kernel の実行は非同期 API を用いることで並列に行うことができる。
+
+通常、フレームワークが依存関係を解決して実行するが、
+非同期 API を用いる場合はユーザが依存関係を考慮する必要がある。
+しかし Task の依存関係は TaskManager が既に解決した状態で送ってくるので、
+ Scheduler は依存関係を考慮せずに実行して問題ない。
+
+GPGPU 用の Scheduler は CommandQueue を2つ持っており、Task をパイプライン的に実行する。
+
+転送されてきた Task が全て終了すると、 
+TaskManager 間の通信を担当する同期キューである mail を通して TaskManager に Task の終了を通知する。
+終了が通知されると TaskManager でその TaskList に関する依存関係が解消され、
+
+GPGPU の Scheduler 内で Platform や Device ID の取得、 Context の生成、 Kernel の Build と Load等も行っており、OD
+並列処理したい計算のみに集中できる。
+
 \section{データ並列}
-\subsection{ベンチマーク}
+並列プログラミングにおいて、明示的な並列化部分はループ部分である。
+ GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。
+プログラムの大部分がループであれば、データ並列による実行だけでプログラムの性能は向上する。
+OpenCL 、 CUDA ともにデータ並列をサポートしている。
+Task を実行する際にデータをどう分割するか指定し、kernel をデータ並列実行用に書き換えることで実現する。
+データ並列実行用の kernel は以下のように記述する。
+2つの input データの積を output データに格納する例題、 multiply を用いる。
+
+\begin{lstlisting}[frame=lrbt,label=src:multiply_opencl,caption=Multiply(OpenCL),numbers=left]
+__kernel void
+multiply(__global const long  *params,
+         __global const float *input1,
+         __global const float *input2,
+         __global const float *output) {
+
+    long id = get_global_id(0);
+
+    output[id] = input1[id] * input2[id];
+}
+\end{lstlisting}
+
+\begin{lstlisting}[frame=lrbt,label=src:multiply_cuda,caption=Multiply(CUDA),numbers=left]
+__global__ void 
+multiply(__global const long  *params,
+         __global const float *input1,
+         __global const float *input2,
+         __global const float *output) {
+
+    int id = blockIdx.x * blockDim.x + threadIdx.x;
+
+    output[id] = input1[id] * input2[id];
+}
+\end{lstlisting}
+
+このような kernel を分割数分生成する。
+分割数は kernel の生成時にそれぞれのフレームワークが 用意している API を用いて指定する。
+いずれの kernel も
+\begin{itemize}
+\item 自分の計算する範囲を取得(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の7行目)
+\item 取得した範囲を計算(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}の9行目)
+\end{itemize}
+という手順で処理する。
+計算する範囲については OpenCL では取得用の API を用い、 CUDA では kernel の持つ組み込み変数から算出する。
+
+Cerium でも データ並列実行をサポートする。
+ GPU におけるデータ並列実行だけでなくマルチコア CPU 上でのデータ並列実行にも対応する。
+なお、マルチコア CPU 上で実行する場合も GPU 実行時の kernel
+(ソースコード\ref{src:multiply_opencl}、ソースコード\ref{src:multiply_cuda}) となるべく近い形式で記述できるようにする。
+マルチコア CPU 上でデータ並列実行する場合、 kernel は以下のように記述する。
+
+\begin{lstlisting}[frame=lrbt,label=src:multiply_cpu, caption=Multiply(CPU),numbers=left]
+static int 
+run(SchedTask *s, void *rbuf, void *wbuf) {
+    float *indata1, *indata2, *outdata;
+
+    indata1 = (float*)s->get_input(rbuf, 0);
+    indata2 = (float*)s->get_input(rbuf, 0);
+    outdata = (float*)s->get_output(wbuf, 0);
+
+    long id = (long)s->get_param(0);
+    outdata[id] = indata1[id] * indata2[id];
+    return 0;
+}
+\end{lstlisting}
+
+OpenCL 、CUDA と違い値を引数として直接渡すのではなく、メモリバッファから Load し、計算を行う。
+値渡しや修飾子等若干の違いはあるが、ほぼ同じ形式で kernel を記述することができる。
+CPU、 OpenCL、 CUDA いずれか1つの記述から残りのコードも生成できるようにする事が望ましい。
+
+Cerium でデータ並列実行を行う場合、Task を spwan API でなく iterate API で生成すればよい。
+iterate API は複数の length を引数とし、
+length の値がデータ分割後に各 Task が担当するサイズ、length の個数がデータの次元数となる。
+これを元にScheduler が各 Task が担当する index を計算し、Task に set\_param する。
+
+Task は実行時に get\_param することで set\_param した値を取得し、担当範囲をデータ並列を実行する。
+この get\_param が OpenCL における get\_global\_id API に相当する。
+
+index の割り当ての例を示す。
+データ数10個の入力を持つ Task に対して CPU 数4、
+一次元における分割でデータ並列実行した場合の index の割り当ては表:\ref{table:dataparallel_index}になる。
+
+この例だと各 CPU に対する index の割り当ては CPU0 は index 0、4、8、 CPU1 は index 1、5、9、
+CPU2 は index 2、6、CPU3 は index 3、7となる。
+
+\begin{tiny}
+  \begin{table}[htpb]
+    \begin{center}
+      \small
+      \begin{tabular}[htpb]{c||c|c|c|c}
+        \hline
+        stage & CPU0 & CPU1 & CPU2 & CPU3 \\
+        \hline
+        \hline
+        1 & 0 & 1 & 2 & 3 \\
+        \hline
+        2 & 4 & 5 & 6 & 7 \\
+        \hline
+        3 & 8 & 9 &   &   \\
+        \hline
+      \end{tabular}
+      \caption{データ並列実行時の index の割り当て}
+      \label{table:dataparallel_index}
+    \end{center}
+  \end{table}
+\end{tiny}
+
+並列プログラミングだと、並列化部分が全て同一の Task であるということは少なくない。
+その際、 Task 生成部分をループで回すことなく、簡単な Syntax で記述することができる。
+
+データ並列で実行する場合、 Input と Output を各 Task 間で共有するため、少ないコピーに抑えられる。
+CPU ではメモリ領域を節約する事ができるが、 Task と Manager でメモリ領域が同じ(\ref{sec:shared_memory}節)なため、
+コピーによるオーバーヘッドは少ない。
+
+しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなる。
+コピーを減らす事で並列度の向上が見込める。