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