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