--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/paper/data_parallel.tex	Mon Nov 04 23:35:10 2013 +0900
@@ -0,0 +1,87 @@
+\section{Ceriumにおけるデータ並列}\label{data_parallel}
+OpenCLで充分な並列度を得るには、データ並列による実行をサポートした方が良い。
+CeriumでOpenCLのデータ並列を使うために、iterateというAPIを用意した。
+
+ベンチマークをとるために、まずはCPU(many core)上でデータ並列の機構を実装した。
+OpenCLでデータ並列を行う際は、NDRangeの引数でワークアイテムのサイズを設定し、以下のようにkernelを書けばよい。
+
+\begin{Verbatim}[fontsize=\footnotesize,xleftmargin=1cm]
+__kernel void
+multi(__global const float *i_data1, 
+      __global const float *i_data2,
+      __global float *o_data)
+{
+    int i = get_global_id(0);
+    o_data[i] = i_data1[i]*i_data2[i];
+}
+
+\end{Verbatim}
+kernelを複数生成し、各kernelは自分が担当するindexをget\_global\_id APIで取得し、
+その部分だけ計算を行う。CPUで実行する場合もGPU実行時のkernelとなるべく近い形式で記述できるようにする。
+
+\subsection{データ並列実行の機構}
+データ並列で実行する場合はspawn APIではなく、iterate APIでTaskを生成すればよい。
+Scheduler内で引数分のTaskを生成し、それぞれに自分が担当するindexをパラメタとして設定していく。
+iterateにはlengthを引数として渡し、lengthの値と渡したlengthの個数で次元数や
+ワークアイテムのサイズをSchedulerが計算する。
+CPU実行時のkernelは以下のように記述する。
+
+\begin{Verbatim}[fontsize=\footnotesize,xleftmargin=1cm]
+static int // kernel
+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);
+
+    long i = (long)s->get_param(0);
+    outdata[i]=indata1[i]*indata2[i];
+    return 0;
+}
+\end{Verbatim}
+
+\subsection{データ並列におけるindex割り当ての実装}
+Taskを生成するとき、dimensionとワークアイテムのサイズをもとに各Taskが担当するindexを計算し、set\_paramする。
+kernelはget\_paramでそのindexを取得してデータ並列で実行する。
+get\_param APIがOpenCLのget\_global\_id APIに相当する。
+
+例として、cpu数4、一次元で10個のdataにたいしてデータ並列実行を行った場合、
+各CPUが担当するindexは表:\ref{table:data_parallel_index}のようになる。
+
+この例だと各CPUに対するindexの割り当ては、
+CPU0はindex0、4、8、
+CPU1はindex1、5、9、
+CPU2はindex2、6、
+CPU3はindex3、7となっている。
+
+\begin{tiny}
+  \begin{table}[h]
+    \begin{center}
+      \caption{data並列実行時のindexの割り当て}
+      \label{table:data_parallel_index}
+      \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}
+  \end{table}
+
+\end{tiny}
+この実装により、Ceriumでデータ並列の実行が可能になった。
+並列プログラミングだと、並列化するTaskが全部同一であるという事は少なくない。
+その際、Taskを生成する部分をループで回すことなく、簡単なsyntaxで記述できる。
+
+データ並列で実行する場合は、inputとoutputを各Taskで共有するため、少ないコピーですむ。
+CPUならメモリ領域がTaskとmanagerで同じなので、dataのコピーで大きいオーバーヘッドにはならない。
+しかしCellとGPUはメモリ領域が異なるため、dataコピーのオーバーヘッドが大きく、
+データ並列による高速化が見込める。