Mercurial > hg > Members > kokubo > 2013-mid-thesis
comparison paper/data_parallel.tex @ 0:9e88a388ec83
first commit
author | Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp> |
---|---|
date | Tue, 05 Nov 2013 23:18:04 +0900 |
parents | |
children | f4b3de446113 |
comparison
equal
deleted
inserted
replaced
-1:000000000000 | 0:9e88a388ec83 |
---|---|
1 \section{Ceriumにおけるデータ並列}\label{data_parallel} | |
2 OpenCLで充分な並列度を得るには、データ並列による実行をサポートした方が良い。 | |
3 CeriumでOpenCLのデータ並列を使うために、iterateというAPIを用意した。 | |
4 | |
5 ベンチマークをとるために、まずはCPU(many core)上でデータ並列の機構を実装した。 | |
6 OpenCLでデータ並列を行う際は、NDRangeの引数でワークアイテムのサイズを設定し、以下のようにkernelを書けばよい。 | |
7 | |
8 \begin{Verbatim}[fontsize=\footnotesize,xleftmargin=1cm] | |
9 __kernel void | |
10 multi(__global const float *i_data1, | |
11 __global const float *i_data2, | |
12 __global float *o_data) | |
13 { | |
14 int i = get_global_id(0); | |
15 o_data[i] = i_data1[i]*i_data2[i]; | |
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の個数で次元数や | |
26 ワークアイテムのサイズをSchedulerが計算する。 | |
27 CPU実行時のkernelは以下のように記述する。 | |
28 | |
29 \begin{Verbatim}[fontsize=\footnotesize,xleftmargin=1cm] | |
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} | |
44 | |
45 \subsection{データ並列におけるindex割り当ての実装} | |
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} | |
78 | |
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 データ並列による高速化が見込める。 |