annotate paper/data_parallel.tex @ 1:764783c403bd

commit
author Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
date Mon, 01 Apr 2013 19:09:39 +0900
parents
children 771aaa69c616
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
1
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
1 \section{Ceriumにおけるデータ並列}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
2 OpenCLで充分な並列度を得るには、データ並列による実行をサポートした方が良い。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
3 ベンチマークをとるために、まずはCPU(many core)上でデータ並列の機構を実装した。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
4 OpenCLでデータ並列を行う際は、NDRangeの引数でワークアイテムのサイズを設定し、以下のようにkernelを書けばよい。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
5
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
6 \begin{verbatim}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
7 __kernel void
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
8 multi(__global const float *A,
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
9 __global const float*B,
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
10 __global float *C)
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
11 {
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
12 int i = get_global_id(0);
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
13 C[i] = A[i]*B[i];
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
14 }
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
15
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
16 \end{verbatim}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
17 kernelを複数生成し、各kernelは自分が担当するindexをget\_global\_id APIで取得し、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
18 その部分だけ計算を行う。CPUで実行する場合もGPU実行時のkernelとなるべく近い形式で記述できるようにする。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
19
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
20 \subsection{データ並列実行の機構}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
21 データ並列で実行する場合はspawn APIではなく、iterate APIでtaskを生成すればよい。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
22 Scheduler内で引数分taskを生成し、それぞれに自分が担当するindexをパラメタとして設定していく。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
23 iterateにはlengthを引数として渡し、lengthの値と渡したlengthの個数でdimensionや
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
24 ワークアイテムのサイズをSchedulerが計算する。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
25 CPU実行時のkernelは以下のように記述する。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
26
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
27 \begin{verbatim}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
28 static int // kernel
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
29 run(SchedTask *s,void *rbuf, void *wbuf)
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
30 {
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
31 float *indata1,*indata2,*outdata;
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
32
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
33 indata1 = (float*)s->get_input(rbuf, 0);
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
34 indata2 = (float*)s->get_input(rbuf, 1);
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
35 outdata = (float*)s->get_output(wbuf, 0);
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
36
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
37 long i = (long)s->get_param(0);
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
38 outdata[i]=indata1[i]*indata2[i];
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
39 return 0;
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
40 }
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
41 \end{verbatim}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
42 taskを生成するとき、dimensionとワークアイテムのサイズをもとに各taskが担当するindexを計算し、set\_paramする。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
43 kernelはget\_paramでそのindexを取得してデータ並列で実行する。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
44 get\_param APIがopenCLのget\_global\_id APIに相当する。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
45
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
46 例として、cpu数4、一次元で10個のdataにたいしてデータ並列実行を行った場合、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
47 各CPUが担当するindexは表:\ref{table:data_parallel_index}のようになる。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
48
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
49 この例だと各CPUに対するindexの割り当ては、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
50 CPU0はindex0、4、8、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
51 CPU1はindex1、5、9、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
52 CPU2はindex2、6、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
53 CPU3はindex3、7となっている。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
54
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
55 \begin{tiny}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
56 \begin{table}[h]
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
57 \begin{center}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
58 \caption{data並列実行時のindexの割り当て}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
59 \label{table:data_parallel_index}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
60 \small
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
61 \begin{tabular}[t]{c||c|c|c|c}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
62 \hline
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
63 stage&CPU0& CPU1&CPU2&CPU3 \\
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
64 \hline
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
65 1&0&1&2&3 \\
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
66 \hline
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
67 2&4&5&6&7 \\
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
68 \hline
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
69 3&8&9& & \\
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
70 \hline
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
71 \end{tabular}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
72 \end{center}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
73 \end{table}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
74 \end{tiny}
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
75 この実装により、Ceriumでデータ並列の実行が可能になった。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
76 並列プログラミングだと、並列化するtaskが全部同一であるという事は少なくない。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
77 その際、task生成部分を何回もループで回すことなく、簡単なsyntaxで記述できる。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
78
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
79 データ並列で実行する場合は、inputとoutputを各taskで共有するため、少ないコピーですむ。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
80 CPUならメモリ領域がtaskとmanagerで同じなので、dataのコピーで大きいオーバーヘッドにはならない。
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
81 しかしCellとGPUはメモリ領域が異なるため、dataコピーのオーバーヘッドが大きく、
Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
parents:
diff changeset
82 データ並列による高速化が見込める。