--- a/paper/chapter5.tex	Sun Feb 15 12:04:36 2015 +0900
+++ b/paper/chapter5.tex	Sun Feb 15 20:12:33 2015 +0900
@@ -40,12 +40,20 @@
 
 \section{データ並列}
 並列プログラミングにおいて、明示的な並列化部分はループ部分である。
- GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。
+GPU は数百個のコアを有しており、ループ部分に対してデータ並列で処理を行うことで CPU より高速に演算を行う事ができる。
 プログラムの大部分がループであれば、データ並列による実行だけでプログラムの性能は向上する。
+
 OpenCL 、 CUDA ともにデータ並列をサポートしている。
-Task を実行する際にデータをどう分割するか指定し、kernel をデータ並列実行用に書き換えることで実現する。
-データ並列実行用の kernel は以下のように記述する。
-2つの input データの積を output データに格納する例題、 multiply を用いる。
+OpenCL と CUDA はTask を実行する際にデータをどう分割するか指定し、
+kernel にデータ並列用の処理を加えることで可能となる。
+\ref{sec:multicore_dataparallel}節で Cerium でマルチコア CPU におけるデータ並列を可能にした。
+GPGPU においてもデータ並列実行をサポートする。
+GPU 上でのデータ並列実行もマルチコア CPU と変わらず、iterate API によりデータ並列用の Task を生成することができる。
+iterate で Task を生成することで Scheduler が OpenCL 及び CUDA の API に適切なパラメタを渡している。
+Task の生成部分は マルチコア CPU と GPU で完全に同じ形式で記述できる。
+
+データ並列実行の際、Task は以下のように記述する。
+なお、例題は multiply を用いている。
 
 \begin{lstlisting}[frame=lrbt,label=src:multiply_opencl,caption=Multiply(OpenCL),numbers=left]
 __kernel void
@@ -73,84 +81,21 @@
 }
 \end{lstlisting}
 
-このような kernel を分割数分生成する。
-分割数は kernel の生成時にそれぞれのフレームワークが 用意している API を用いて指定する。
-いずれの kernel も
+このような Task を分割数分生成する。
+分割数は Task それぞれのフレームワークが用意している API を用いて指定する。
+
 \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 を記述することができる。
+いずれの Task も上記の手順で処理を行っている。
+計算する範囲について、 OpenCL では取得用の API を用い、 CUDA では kernel の持つ組み込み変数から算出する。
+マルチコア CPU では引数としてデータを直接渡していたが、OpenCL 、CUDA では上記の方法でメモリバッファから Load し、計算を行う。
+値渡しや修飾子等若干の違いはあるが、OpenCL 、CUDA ともにマルチコア CPU(ソースコード:\ref{src:multicore_cpu}) とほぼ同じ形式で 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}節)なため、
+CPU ではメモリ領域を節約する事はできるが、 Task と Manager でメモリ領域が同じ(\ref{sec:shared_memory}節)なため、
 コピーによるオーバーヘッドは少ない。
 
-しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなる。
-コピーを減らす事で並列度の向上が見込める。
+しかし GPU は SharedMemory ではなく、データの転送がオーバーヘッドとなるため、コピーを減らす事で並列度の向上が見込める。