|
58
|
1 <!DOCTYPE html>
|
|
|
2 <html>
|
|
|
3 <head>
|
|
|
4 <meta charset='utf-8'>
|
|
|
5 <title>Seminar</title>
|
|
|
6
|
|
|
7 <!--
|
|
|
8 Notes on CSS media types used:
|
|
|
9
|
|
|
10 1) projection -> slideshow mode (display one slide at-a-time; hide all others)
|
|
|
11 2) screen -> outline mode (display all slides-at-once on screen)
|
|
|
12 3) print -> print (and print preview)
|
|
|
13
|
|
|
14 Note: toggle between projection/screen (that is, slideshow/outline) mode using t-key
|
|
|
15
|
|
|
16 Questions, comments?
|
|
|
17 - send them along to the mailinglist/forum online @ http://groups.google.com/group/webslideshow
|
|
|
18 -->
|
|
|
19
|
|
|
20 <!-- style sheet links -->
|
|
|
21 <link rel="stylesheet/less" href="themes/blank/projection.css.less" media="screen,projection">
|
|
|
22 <link rel="stylesheet/less" href="themes/blank/screen.css.less" media="screen">
|
|
|
23 <link rel="stylesheet/less" href="themes/blank/print.css.less" media="print">
|
|
|
24
|
|
|
25 <link rel="stylesheet/less" href="blank.css.less" media="screen,projection">
|
|
|
26
|
|
|
27 <!-- Notes about less css support
|
|
|
28 - all less stylesheets (*.css.less) need to get listed/loaded first (before the less.js script)
|
|
|
29 - find more info about less.js online @ http://lesscss.org
|
|
|
30
|
|
|
31 ***** NOTE:
|
|
|
32 less.js browser script currently won’t work if you’re using Google Chrome
|
|
|
33 and the path to your page starts with "file:///" due to a known Chrome issue.
|
|
|
34 (In the developer/js console you will see:
|
|
|
35 XMLHttpRequest cannot load file:///../s6/shared/projection.css.less.
|
|
|
36 Cross origin requests are only supported for HTTP.)
|
|
|
37 -->
|
|
|
38
|
|
|
39 <!-- add js libs (less, jquery) -->
|
|
|
40 <script src="js/less-1.1.4.min.js"></script>
|
|
|
41 <script src="js/jquery-1.7.min.js"></script>
|
|
|
42
|
|
|
43 <!-- S6 JS -->
|
|
|
44 <script src="js/jquery.slideshow.js"></script>
|
|
|
45 <script src="js/jquery.slideshow.counter.js"></script>
|
|
|
46 <script src="js/jquery.slideshow.controls.js"></script>
|
|
|
47 <script src="js/jquery.slideshow.footer.js"></script>
|
|
|
48 <script src="js/jquery.slideshow.autoplay.js"></script>
|
|
|
49 <script>
|
|
|
50 $(document).ready( function() {
|
|
|
51 Slideshow.init();
|
|
|
52
|
|
|
53 // Example 2: Start Off in Outline Mode
|
|
|
54 // Slideshow.init( { mode: 'outline' } );
|
|
|
55
|
|
|
56 // Example 3: Use Custom Transition
|
|
|
57 // Slideshow.transition = transitionScrollUp;
|
|
|
58 // Slideshow.init();
|
|
|
59
|
|
|
60 // Example 4: Start Off in Autoplay Mode with Custom Transition
|
|
|
61 // Slideshow.transition = transitionScrollUp;
|
|
|
62 // Slideshow.init( { mode: 'autoplay' } );
|
|
|
63 } );
|
|
|
64 </script>
|
|
|
65
|
|
|
66 <!-- Better Browser Banner for Microsoft Internet Explorer (IE) -->
|
|
|
67 <!--[if IE]>
|
|
|
68 <script src="js/jquery.microsoft.js"></script>
|
|
|
69 <![endif]-->
|
|
|
70
|
|
|
71 </head>
|
|
|
72 <body>
|
|
|
73
|
|
|
74 <div class="layout">
|
|
|
75 <div id="header"></div>
|
|
|
76 <div id="footer">
|
|
|
77 <div align="right">
|
|
|
78 <img src="images/concurrency.png" width="200">
|
|
|
79 </div>
|
|
|
80 </div>
|
|
|
81 </div>
|
|
|
82
|
|
|
83 <div class="presentation">
|
|
|
84
|
|
|
85 <!-- add slides here; example -->
|
|
|
86
|
|
|
87 <div class='slide cover'>
|
|
|
88 <table width="90%" height="90%" border="0" align="center">
|
|
|
89 <tr>
|
|
|
90 <td><div align="center">
|
|
|
91 <h1><font color="#808db5">マルチプラットフォーム対応<br>並列プログラミングフレームワーク</font></h1>
|
|
|
92 </div></td>
|
|
|
93 </tr>
|
|
|
94 <tr>
|
|
|
95 <td><div align="left">
|
|
|
96 Yuhi TOMARI
|
|
|
97 <script>
|
|
|
98 var date = new Date();
|
|
|
99 var year = date.getFullYear();
|
|
|
100 var month = date.getMonth();
|
|
|
101 var day = date.getDate();
|
|
|
102
|
|
|
103 var monthList = new Array("January","February","March","April","May","June",
|
|
|
104 "July","August","September","October","November","December");
|
|
|
105
|
|
|
106 document.write(monthList[month]+" "+day+", "+year);
|
|
|
107
|
|
|
108 </script>
|
|
|
109 <hr style="color:#ffcc00;background-color:#ffcc00;text-align:left;border:none;width:300%;height:0.2em;">
|
|
|
110 </div></td>
|
|
|
111 </tr>
|
|
|
112 </table>
|
|
|
113 </div>
|
|
|
114
|
|
|
115 <div class='slide'>
|
|
61
|
116 <h2>マルチプラットフォームなフレームワークにおける並列プログラミング 1/2</h2>
|
|
63
|
117 <p>
|
|
|
118 消費電力や発熱、クロックの限界といった問題から CPU の性能を上げることによる処理性能の向上は難しい。
|
|
|
119 マルチコア CPU や GPU を含んだヘテロジニアス構成が主流になっている。
|
|
58
|
120 クロックの性能を上げるのではなく、コア数を増やす事でパフォーマンスを向上させている。
|
|
|
121 </p>
|
|
|
122 <p>
|
|
|
123 マルチコア CPU や GPU といった<font color="red">マルチコアプラットフォーム</font>なアーキテクチャ上で
|
|
|
124 リソースを有効活用するには、それぞれのプラットフォームに最適な形でプログラムを並列に動作させる必要がある。
|
|
|
125 </p>
|
|
|
126 <p>しかしこれらのチューニングは複雑で、コーディング時に毎回行うと複雑さや拡張性の問題がある。</p>
|
|
|
127 </div>
|
|
|
128
|
|
|
129
|
|
|
130 <div class='slide'>
|
|
61
|
131 <h2>マルチプラットフォームなフレームワークにおける並列プログラミング 2/2</h2>
|
|
58
|
132 <p>
|
|
|
133 そういった問題を解決するため、本研究では並列プログラミングフレームワーク、 Cerium の開発を行った。
|
|
|
134 異なるプラットフォーム上で最適なチューニングを行うため、以下の実装を行った。
|
|
|
135 </p>
|
|
|
136 <ul>
|
|
|
137 <li>パイプライニングによる Task の並列実行</li>
|
|
|
138 <li>OpenCL、CUDA を用いた GPGPU 対応</li>
|
|
|
139 <li>データ並列実行</li>
|
|
|
140 <li>並列処理むけのI/O</li>
|
|
|
141 </ul>
|
|
|
142 <p>
|
|
|
143 Sort、WordCount、FFT といった例題を元に、これら Cerium の並列実行機構が
|
|
|
144 マルチプラットフォームにおける並列プログラミングで有効に作用することを示す。
|
|
|
145 </p>
|
|
|
146 </div>
|
|
|
147
|
|
|
148 <div class='slide'>
|
|
|
149 <h2>並列プログラミングフレームワーク Cerium</h2>
|
|
|
150 <p>
|
|
|
151 Cerium は Linux、MacOSX 上で動作する汎用計算用の並列プログラミングフレームワークである。
|
|
|
152 </p>
|
|
63
|
153 <div align="center">
|
|
|
154 <img src="./images/cerium_image.png" width="700">
|
|
|
155 </div>
|
|
58
|
156 <p>Cerium を用いることでマルチコア CPU と GPU において Scheduling を含めたプログラミングを可能となる。</p>
|
|
|
157 </div>
|
|
|
158
|
|
|
159 <div class='slide'>
|
|
60
|
160 <h2>Cerium における Task の生成</h2>
|
|
58
|
161
|
|
|
162 <p>
|
|
|
163 Cerium TaskManager では処理の単位を Task としてプログラムを記述していく。
|
|
|
164 関数やサブルーチンを Task として扱い、Task に各種パラメタを設定した後に並列実行される。
|
|
|
165 Input データの各要素同士を乗算し、 Output に格納する Multiply という例題がある。
|
|
|
166 Multiply の例題を元に Cerium で Task が生成される様子を以下に示す。
|
|
|
167 </p>
|
|
|
168 <pre class="code">
|
|
|
169 void
|
|
|
170 multiply_init(TaskManager *manager, float *i_data1,
|
|
|
171 float *i_data2, float *o_data) {
|
|
|
172
|
|
|
173 // create task
|
|
|
174 HTask* multiply = manager->create_task(MULTIPLY_TASK);
|
|
|
175 multiply->set_cpu(spe_cpu);
|
|
|
176
|
|
|
177 // set indata
|
|
|
178 multiply->set_inData(0, i_data1, sizeof(float) * length);
|
|
|
179 multiply->set_inData(1, i_data2, sizeof(float) * length);
|
|
|
180
|
|
|
181 // set outdata
|
|
|
182 multiply->set_outData(0, o_data, sizeof(float) * length);
|
|
|
183
|
|
|
184 // set parameter
|
|
|
185 multiply−>set_param(0,(long)length);
|
|
|
186
|
|
|
187 // set device
|
|
|
188 multiply->set_cpu(SPE_ANY);
|
|
|
189
|
|
|
190 // spawn task
|
|
|
191 multiply−>spawn();
|
|
|
192 }
|
|
|
193 </pre>
|
|
|
194 </div>
|
|
|
195
|
|
|
196 <div class='slide'>
|
|
60
|
197 <h2>Cerium における Task の記述</h2>
|
|
|
198 <p>Device 側で実行される Task の記述を示す。</p>
|
|
|
199 <pre class="code">
|
|
|
200 static int
|
|
|
201 run(SchedTask ∗s) {
|
|
|
202 float ∗i_data1 = (float∗)s−>get_input(0); // get input
|
|
|
203 float ∗i_data2 = (float∗)s−>get_input(1); // get output
|
|
|
204 float ∗o_data = (float∗)s−>get_output(0); // get parameter
|
|
|
205 long length = (long)s−>get_param(0);
|
|
|
206
|
|
|
207 // calculate
|
|
|
208 for (int i=0; i<length; i++) {
|
|
|
209 o_data[i] = i_data1[i] ∗ i_data2[i];
|
|
|
210 }
|
|
|
211 return 0;
|
|
|
212 }
|
|
|
213 </pre>
|
|
|
214 <p>Host 側では Task を生成する際に様々なパラメタを設定しており、
|
|
|
215 Task にはそれを取得する API が用意されている。</p>
|
|
|
216 <table border="0" >
|
|
|
217 <tr bgcolor="palegreen">
|
|
|
218 <th align="center">API</th><th align="center">content</th>
|
|
|
219 </tr>
|
|
|
220
|
|
|
221 <tr bgcolor="dbffa3">
|
|
|
222 <th align="left" >get_input</th><th align="left">入力データのアドレスを取得</th>
|
|
|
223 </tr>
|
|
|
224 <tr bgcolor="palegreen">
|
|
|
225 <th align="left">get_output</th><th align="left">出力先データのアドレスを取得</th>
|
|
|
226 </tr>
|
|
|
227 <tr bgcolor="dbffa3">
|
|
|
228 <th align="left">get_param</th><th align="left">パラメータを取得</th>
|
|
|
229 </tr>
|
|
|
230 </table>
|
|
|
231 </div>
|
|
|
232
|
|
|
233 <div class='slide'>
|
|
63
|
234 <h2>Task の依存関係の記述</h2>
|
|
|
235 <p>
|
|
|
236 並列処理を行う場合、Task を大量に生成する場合がある。
|
|
|
237 そういった場合において一括で Task を生成/実行してしまうと並列度が落ちてしまう。
|
|
|
238 これは生成しただけで Task そのものがメモリを圧迫してしまっていることが原因となる。
|
|
|
239 </p>
|
|
|
240 <p>
|
|
|
241 そういった 例題において、Task は一定数ずつ徐々に生成/実行する必要がある。
|
|
|
242 ということは、Block 間で依存関係を設定する必要がある。
|
|
|
243 依存関係について Cerium の Bitonic Sort を例題に考える。
|
|
|
244 </p>
|
|
|
245 </div>
|
|
|
246
|
|
|
247
|
|
|
248 <div class='slide'>
|
|
|
249 <h2>Bitonic Sort の例題</h2>
|
|
|
250 <p>Bitonic Sort は配列の分割を行い、分割した部分に対して Sort を行う。
|
|
|
251 分割後の Sort には QuickSort を使用している。Bitonic Sort は2つの Sort を行う。
|
|
|
252 </p>
|
|
|
253 <ul>
|
|
|
254 <li>使用する CPU 数を元に分割数を算出し、分割した箇所に対して Sort する(fsort)
|
|
|
255 <li>Block の中間から次の Block の中間までを Sort する(bsort)
|
|
|
256 </ul>
|
|
|
257 <p>この2つの Sort を分割数分繰り返している</p>
|
|
|
258 </div>
|
|
|
259
|
|
|
260 <div class='slide'>
|
|
|
261 <h2>Bitonic Sort の例題</h2>
|
|
|
262 <div align="center">
|
|
|
263 <img src="./images/fsort_bsort.png" width="850">
|
|
|
264 </div>
|
|
|
265 </div>
|
|
|
266
|
|
|
267 <div class='slide'>
|
|
|
268 <h2>Task 間の依存関係</h2>
|
|
|
269 <p>Bitonic Sort を行う際、依存関係として bsort は fsort の結果に対して sort を行い、
|
|
|
270 fsort は前の Stage の bsort に対して Sort を行う必要がある
|
|
|
271 </p>
|
|
|
272 <p>よって、BitonicSort のような大量に Task を生成する例題を並列実行する場合、
|
|
|
273 「例題の性質としての依存関係」と「Task を徐々に生成するための依存関係」
|
|
|
274 の二種類の依存関係を記述する必要がある。</p>
|
|
|
275 </div>
|
|
|
276
|
|
|
277 <div class='slide'>
|
|
|
278 <h2>依存関係の記述</h2>
|
|
|
279 <p>例題独自の依存関係</p>
|
|
|
280 <pre class="code" align="left">static int
|
|
|
281 sort_start(SchedTask *manager, void *d, void *e)
|
|
|
282 {
|
|
|
283 Sort *s = (Sort*)manager->get_param(0);
|
|
|
284 long half_num = s->split_num-1;
|
|
|
285
|
|
|
286 for (int i = 0; i < s->split_num-1; i++) {
|
|
|
287 s->fsort[i] = manager->create_task(QUICK_SORT,(memaddr)&s->data[i*block_num],
|
|
|
288 sizeof(Data)*block_num,
|
|
|
289 (memaddr)&s->data[i*block_num],
|
|
|
290 sizeof(Data)*block_num);
|
|
|
291
|
|
|
292 s->fsort[i]->wait_for(s->bsort[i-1]);
|
|
|
293 }
|
|
|
294 ~省略~
|
|
|
295 </pre>
|
|
|
296 </div>
|
|
|
297 <div class='slide'>
|
|
|
298 <h2>依存関係の記述</h2>
|
|
|
299 <p>Task を徐々に生成するための依存関係</p>
|
|
|
300 <pre class="code" align="left">
|
|
|
301 // recursive Task
|
|
|
302 HTaskPtr restart = manager->create_task(SortSimple,0,0,0,0);
|
|
|
303 restart->set_param(0,(memaddr)s);
|
|
|
304 restart->wait_for(s->fsort[0]);
|
|
|
305 for (int i = 0; i < s->split_num; i++) {
|
|
|
306 s->fsort[i]->spawn();
|
|
|
307 }
|
|
|
308 restart->spawn();
|
|
|
309 return 0;
|
|
|
310 }
|
|
|
311 </pre>
|
|
|
312 </div>
|
|
|
313
|
|
|
314 <div class='slide'>
|
|
60
|
315 <h2>TaskManager の構成</h2>
|
|
|
316 <div align="center">
|
|
|
317 <img src='images/createtask.png' width="700">
|
|
|
318 </div>
|
|
|
319 <ul>
|
|
|
320 <li>TaskManagerと各Threadsの間には Syncronized な Mail Queueがある。
|
|
|
321 <li>依存関係の解決された Task は TaskManager から Mail Queue に送られる。
|
|
|
322 <li>Task に設定された CPUType に対応した Threads が Mail Queue から Task を取得し、並列実行していく。
|
|
|
323 </ul>
|
|
|
324 </div>
|
|
|
325
|
|
|
326 <div class='slide'>
|
|
|
327 <h2>マルチコア CPU 上での並列実行</h2>
|
|
|
328 <div align="center">
|
|
|
329 <img src="images/pipeline.png" width="600">
|
|
|
330 </div>
|
|
|
331 <p>
|
|
|
332 Cerium は Cell 上で動作するフレームワークであったが MacOSX、Linux 上での並列実行に対応させた。
|
|
|
333 </p>
|
|
|
334 <p>
|
|
|
335 マルチコア CPU 上での並列実行は、Synchronized Queue とパイプラインによって実現されている。
|
|
|
336 TaskManager で依存関係を解決された Task は Scheduler に送信され、
|
|
|
337 Scheduler が持っているパイプラインの機構に沿って並列に実行する。
|
|
|
338 </p>
|
|
|
339 </div>
|
|
|
340
|
|
|
341 <div class='slide'>
|
|
|
342 <h2>マルチコア CPU におけるパイプラインの実装</h2>
|
|
|
343 <table>
|
|
|
344 <tr>
|
|
|
345 <th>
|
|
|
346 <pre class="code" align="left">void
|
|
|
347 Scheduler::run(SchedTaskBase* task1) {
|
|
|
348 SchedTaskBase* task2 = new SchedNop();
|
|
|
349 SchedTaskBase* task3 = new SchedNop();
|
|
|
350
|
|
|
351 // main loop
|
|
|
352 do {
|
|
|
353
|
|
|
354 task1->read();
|
|
|
355 task2->exec();
|
|
|
356 task3->write();
|
|
|
357
|
|
|
358 delete task3;
|
|
|
359
|
|
|
360 task3 = task2;
|
|
|
361 task2 = task1;
|
|
|
362 task1 = task1->next(this, 0);
|
|
|
363
|
|
|
364 } while (task1);
|
|
|
365
|
|
|
366 delete task3;
|
|
|
367 delete task2;
|
|
|
368 }</pre>
|
|
|
369 </th>
|
|
|
370 <th align="left">
|
|
|
371 <p>
|
|
|
372 Cerium の Task は SchedTask と呼ばれるデータ構造で表現されている。
|
|
|
373 SchedTask は read/exec/write のメソッドを持っており、
|
|
|
374 パイプラインの各ステージで段階的に実行される。
|
|
|
375 </p>
|
|
|
376 <p>
|
|
|
377 引数として TaskList を受け取り、List 内の Task をパイプライン実行する。
|
|
|
378 task3 が write を担当しており、write が終わった Task は終了となる。
|
|
|
379 </p>
|
|
|
380 <p>
|
|
|
381 終了した task は delete して良い。
|
|
|
382 task3=task2、task2=task1 と SchedTask をずらして行き、TaskList から 次の Task を読み込む。
|
|
|
383 </p>
|
|
|
384 </th>
|
|
|
385 </tr>
|
|
|
386 </table>
|
|
|
387 </div>
|
|
|
388
|
|
|
389 <div class='slide'>
|
|
63
|
390 <h2>マルチコア CPU におけるデータ並列</h2>
|
|
|
391 <p>
|
|
|
392 Cerium はタスク並列による実行のみを行っていた。
|
|
|
393 並列化を行う問題によってはデータ並列を行った方が良い場合がある。
|
|
|
394 </p>
|
|
|
395 <p>
|
|
|
396 タスク並列は1つのデータに対して異なる処理方法を適用し、それぞれ独立して実行させるものである。
|
|
|
397 </p>
|
|
|
398
|
|
|
399 <p>
|
|
|
400 データ並列は多くのデータを1つのタスクに与え、データごとに独立した処理を行わせる手法である。
|
|
|
401 </p>
|
|
|
402 <p>処理対象となるデータが充分な数のサブセットへ分割可能な場合、データ並列が有効となる。</p>
|
|
60
|
403 </div>
|
|
|
404
|
|
|
405 <div class='slide'>
|
|
63
|
406 <h2>iterate API</h2>
|
|
|
407 <p>
|
|
|
408 データ並列による実行を行う場合、一つの記述から複数のTaskを生成する必要がある。
|
|
|
409 生成した各TaskにIDとinput/output dataを割り当てる「iterate」というAPIを実装した。
|
|
|
410 </p>
|
|
|
411
|
|
|
412 <table>
|
|
|
413 <tr>
|
|
|
414 <td>
|
|
|
415 <img src="images/iterate.png" height="450"></img>
|
|
|
416 </td>
|
|
|
417 <td>
|
|
|
418 <ul>
|
|
|
419 <li>1つの記述から複数のTaskを生成する</li>
|
|
|
420 <li>生成した複数のTaskにIDとInput/Output Dataを割り当てる</li>
|
|
|
421 </ul>
|
|
|
422 この例だと、Taskの持つidとTaskに割り当てられるデータは
|
|
|
423 1対1で対応している。id=割り当てられたdataのindexとなっている。<br><br>
|
|
|
424 並列プログラミングだと、並列化部分が全て同一の Task であるという場合は少なくない。
|
|
|
425 iterate API ならループで回すような処理をする必要が無く、容易な Syntax で記述できる。
|
|
|
426 </td>
|
|
|
427 </tr></table>
|
|
|
428 </div>
|
|
|
429
|
|
|
430 <div class='slide'>
|
|
|
431 <h2>マルチコア CPU によるデータ並列実行</h2>
|
|
|
432 <p>
|
|
|
433 マルチコア CPU においてデータ並列実行する場合、以下のように記述する。
|
|
|
434 例題として 2つの input のデータの積を output データに格納して返す例題、multiply を用いた。
|
|
|
435 </p>
|
|
|
436 <pre class="code">
|
|
|
437 static int
|
|
|
438 run(SchedTask *s, void *rbuf, void *wbuf) {
|
|
|
439 float *indata1, *indata2, *outdata;
|
|
|
440
|
|
|
441 indata1 = (float*)s->get_input(rbuf, 0);
|
|
|
442 indata2 = (float*)s->get_input(rbuf, 0);
|
|
|
443 outdata = (float*)s->get_output(wbuf, 0);
|
|
|
444
|
|
|
445 long id = (long)s->get_param(0);
|
|
|
446 outdata[id] = indata1[id] * indata2[id];
|
|
|
447 return 0;
|
|
|
448 }
|
|
|
449 </pre>
|
|
|
450 <p>get_param によって自分の担当する index を取得し、担当範囲のみを計算する。</p>
|
|
|
451 <p>データ並列実行する場合、各Task に Input/Outpu を設定するのではなく、
|
|
|
452 全ての Task でデータを共有する。共有したデータの自分の担当する箇所にのみ計算を行う。
|
|
|
453 そのため少ないコピーにおさえることができる。
|
|
|
454 </p>
|
|
|
455 </div>
|
|
|
456
|
|
|
457
|
|
|
458 <div class='slide'>
|
|
|
459 <h2>DMA 転送</h2>
|
|
|
460 <p>Cerium は DMA 転送をサポートしている。
|
|
|
461 DMA とは CPU を介さずに周辺装置とメモリ間でデータ転送を行う転送方式である。
|
|
|
462 </p>
|
|
|
463 <p>
|
|
|
464 DMA は prefetch と呼ばれる転送先読みの機能がある。
|
|
|
465 DMA の転送効率を向上させるために送信データを予め取り込んでおく機能である。
|
|
|
466 prefetch による転送機能を追加した。
|
|
|
467 </p>
|
|
60
|
468 </div>
|
|
|
469
|
|
|
470 <div class='slide'>
|
|
61
|
471 <h2>GPU 上での並列実行</h2>
|
|
63
|
472 <p>
|
|
|
473 GPU 上での並列実行をサポートするフレームワークとして、OpenCL と CUDA が挙げられる。
|
|
|
474 これらのフレームワークを用いて Cerium に GPU 上で 並列実行する機能を加えた。
|
|
|
475 </p>
|
|
|
476 <p>
|
|
64
|
477 TaskManager から受け取った Task やデータをOpenCL、CUDA の API を介して GPU に転送する機構、
|
|
63
|
478 GpuScheduler と CudaScheduler を実装した。
|
|
|
479 </p>
|
|
|
480 <div align="center">
|
|
|
481 <img src="./images/gpu_image.png" width="600">
|
|
|
482 </div>
|
|
60
|
483 </div>
|
|
|
484
|
|
|
485 <div class='slide'>
|
|
63
|
486 <h2>フレームワークを用いた GPU の制御</h2>
|
|
64
|
487 <p>
|
|
|
488 GpuScheduler、CudaScheduler ではそれぞれのフレームワークを用いて GPU の制御を行っている。
|
|
|
489 行われていることは以下の3つに分けられる。
|
|
|
490 </p>
|
|
|
491 <ul>
|
|
|
492 <li>Host から Device へのデータ転送
|
|
|
493 <li>kernel の実行
|
|
|
494 <li>Device から Host へのデータ転送
|
|
|
495 </ul>
|
|
|
496 <p>
|
|
|
497 CommandQueue と呼ばれる機構を用いてこういった GPU を制御するための処理を行っていく。
|
|
|
498 CommandQueue に命令を起こるためのしくみで、制御は全てこの Queue を介して行われる。
|
|
|
499 </p>
|
|
|
500 <p>これらはRead, Exec、Write に対応する。
|
|
|
501 GPGPU 用の Scheduler でもパイプラインを構成する。</p>
|
|
63
|
502 </div>
|
|
|
503
|
|
|
504 <div class='slide'>
|
|
64
|
505 <h2>GPGPU におけるパイプラインの実装(Read)</h2>
|
|
|
506 <p>
|
|
|
507 GpuScheduler では SchedTask を用いてない。
|
|
|
508 メインループでは2つの CommandQueue を保持し、GPU の制御命令を二段のパイプラインで実行していく。
|
|
|
509 TaskList から Task を取り出し、Task から実行する kernel やパラメタを生成し、
|
|
|
510 各種フレームワークの API を通して GPU のメモリに転送される。
|
|
|
511 </p>
|
|
|
512
|
|
|
513 <p>
|
|
|
514 全ての Task が終了すると SynchronizedQueue を通してTaskManager に終了を通知する。
|
|
|
515 </p>
|
|
|
516
|
|
|
517 <p>
|
|
|
518 Scheduler の内部で Platform や DeviceID の取得、
|
|
|
519 kernel の build や load といった API を使用するための初期化も行っており、
|
|
|
520 並列化したい処理のみに集中できる。
|
|
|
521 </p>
|
|
60
|
522 </div>
|
|
|
523
|
|
|
524 <div class='slide'>
|
|
|
525 <h2>GPGPU におけるデータ並列</h2>
|
|
64
|
526 <p>
|
|
|
527 マルチコア CPU と同様に、GPGPU に関してもデータ並列実行をサポートした。
|
|
|
528 </p>
|
|
|
529 <p>
|
|
|
530 GPU 上でデータ並列実行する際も iterate API によりデータ並列用の Task を生成できる。
|
|
|
531 生成部分の記述はマルチコア CPU と同じ形式で記述できる。
|
|
|
532 また、Task 自体の記述もほぼ同じ形式となる。以下に Task の例を示す。
|
|
|
533 </p>
|
|
|
534
|
|
|
535 <table>
|
|
|
536 <tr align="left">
|
|
|
537 <th>
|
|
|
538 <pre class="code">
|
|
|
539 __kernel void // OpenCL
|
|
|
540 multiply(__global const long *params,
|
|
|
541 __global const float *input1,
|
|
|
542 __global const float *input2,
|
|
|
543 __global const float *output) {
|
|
|
544
|
|
|
545 long id = get_global_id(0);
|
|
|
546
|
|
|
547 output[id] = input1[id] * input2[id];
|
|
|
548 }
|
|
|
549 </pre>
|
|
|
550 </th>
|
|
|
551 <th>
|
|
|
552 <pre class="code">
|
|
|
553 __global__ void // CUDA
|
|
|
554 multiply(__global const long *params,
|
|
|
555 __global const float *input1,
|
|
|
556 __global const float *input2,
|
|
|
557 __global const float *output) {
|
|
|
558
|
|
|
559 int id = blockIdx.x * blockDim.x + threadIdx.x;
|
|
|
560
|
|
|
561 output[id] = input1[id] * input2[id];
|
|
|
562 }
|
|
|
563 </pre>
|
|
|
564 </th>
|
|
|
565 </tr>
|
|
|
566 </table>
|
|
|
567 </div>
|
|
|
568
|
|
|
569 <div class='slide'>
|
|
|
570 <h2>並列処理向け I/O</h2>
|
|
|
571 <p>
|
|
|
572 ファイルの読み込みなどの I/O を含むプログラムは、
|
|
|
573 読み込み時間が Task のと比較してオーバーヘッドになることが多い。
|
|
|
574 プログラムの並列化を行ったとしても I/O がボトルネックになってしまうと処理は高速にならない。
|
|
|
575 </p>
|
|
|
576 <p>並列計算と同時に動作する、並列 I/O の実装を行った。</p>
|
|
60
|
577 </div>
|
|
|
578
|
|
|
579 <div class='slide'>
|
|
|
580 <h2>Cerium の I/O(mmap による読み込み)</h2>
|
|
64
|
581 <p>
|
|
|
582 Cerium ではファイルの読み込みを mmap で行っていた。</p>
|
|
|
583 <ul>
|
|
|
584 <li>mmap はまず仮想メモリにファイルをマッピングする。
|
|
|
585 <li>マッピングしたメモリ空間にアクセスがあったら OS が読み込みを行う。
|
|
|
586 <li>mmap は並列に動作せず、逐次処理
|
|
|
587 <li>読み込みが OS 依存となり、環境に左右されやすい
|
|
|
588 <p>並列に動作する I/O の機構が必要である</p>
|
|
|
589 </div>
|
|
|
590
|
|
|
591
|
|
|
592 <div class='slide'>
|
|
|
593 <h2>WordCount</h2>
|
|
|
594 <p>サイズの大きいファイルを読み込む例題、WordCount を元に並列 I/O について考える。</p>
|
|
|
595 <p>
|
|
|
596 WordCount は Input としてファイルを受け取り、ファイルの単語数と行数を集計して表示する例題である。
|
|
|
597 </p>
|
|
|
598 <table>
|
|
|
599 <tr><th><img src="./images/wordcount.png" width="600"></th>
|
|
|
600 <th align="left">
|
|
|
601 <ul>
|
|
|
602 <li>input ファイルを一定の大きさ分割する
|
|
|
603 <li>読み込んだテキストファイルに対してそれぞれ並列に計算を行う
|
|
|
604 <li>PrintTask が計算結果を集計して出力する
|
|
|
605 </ul>
|
|
|
606 </th>
|
|
|
607 </tr>
|
|
|
608 </table>
|
|
|
609 </div>
|
|
|
610
|
|
|
611 <div class='slide'>
|
|
|
612 <h2>BlockedRead による I/O の並列化</h2>
|
|
|
613 <p>ファイルを読み込んで、読み込んだファイルに対して並列実行を行う場合、ファイルを分割して処理を行う。</p>
|
|
|
614 <p>よって読み込みの処理自体を分割し、ある程度の大きさ(Block)ごとに読み込みと Task の実行を行う。</p>
|
|
|
615 <p>読み込みの処理自体を分割して行う。これを BlockedRead と呼ぶ。</p>
|
|
|
616 </p>
|
|
60
|
617 </div>
|
|
|
618
|
|
|
619 <div class='slide'>
|
|
64
|
620 <h2>BlockedRead を用いた WordCount</h2>
|
|
|
621 <div align="center">
|
|
|
622 <img src="./images/blockedread.png" width="600">
|
|
|
623 </div>
|
|
|
624 <p>
|
|
|
625 BlockedRead を用いて WordCount を行う際、読み込み用の Task と
|
|
|
626 読み込んだファイルに対して処理を行う Task の2つを生成する。
|
|
|
627 </p>
|
|
|
628 <p>ファイルを分割して読み込み、
|
|
|
629 読み込んだファイルに対して WordCount を行う一定数のTask(BlockedTask)を割り当てる。
|
|
|
630 Task には依存関係を設定する必要があり、図のTask n+1 はTask nを待つ必要がある。
|
|
|
631 </p>
|
|
|
632 <p>まだ読み込みが終了していない領域に割り当てられた Task が起動してしまう事を防ぐためである。</p>
|
|
|
633 <p>この wait によるロックはオーバーヘッドとなるため、なるべく発生しないことが望ましい。</p>
|
|
60
|
634 </div>
|
|
|
635
|
|
|
636 <div class='slide'>
|
|
|
637 <h2>I/O 専用のThread</h2>
|
|
64
|
638 <p>
|
|
|
639 BlockedRead の依存関係による wait はなるべく発生しないことが望ましい。
|
|
|
640 そのため、BlockedRead は連続で Task の起動を行う必要がある。
|
|
|
641 </p>
|
|
|
642 <p>
|
|
|
643 Cerium には SPE_ANY という Thread があり、この Thread で Task の実行を行うと自動で実行するコアを割り振る。
|
|
|
644 しかし、SPE_ANY で BlockedRead を実行すると BlockedRead 間に別の Task が割り込んでしまう場合がある。
|
|
|
645 </p>
|
|
|
646 <div align="center">
|
|
|
647 <img src="./images/speblockedread.png" width="700">
|
|
|
648 </div>
|
|
|
649 <p>TaskBlock の依存関係によっては wait がかかってしまう。そこで、I/O 専用の Thread を作成した。</p>
|
|
60
|
650 </div>
|
|
|
651
|
|
|
652 <div class='slide'>
|
|
64
|
653 <h2>I/O 専用のThread</h2>
|
|
|
654 <p>
|
|
|
655 IO 専用の Thread を作成したが、それだけでは問題は解決しない場合がある。
|
|
|
656 IO thread 内では割り込みが生じる可能性はないが、thread レベルで割り込みが起きる可能性がある。
|
|
|
657 IO thread-SPE_ANY-IO Thread のような実行順序となる場合である。
|
|
|
658 </p>
|
|
|
659 <div align="center">
|
|
|
660 <img src="./images/iothread.png" width="700">
|
|
|
661 </div>
|
|
|
662 <p>
|
|
|
663 そのため、pthread_getschedparam() という POSIX スレッドの API を用いて IO Thread の priority を高く設定した。
|
|
|
664 IO Thread は必ず連続で行われることになる。
|
|
|
665 </p>
|
|
|
666 </div>
|
|
|
667
|
|
|
668 <div class='slide'>
|
|
|
669 <h2>-</h2>
|
|
60
|
670 </div>
|
|
|
671
|
|
|
672 <div class='slide'>
|
|
|
673 <h2>実験に利用する例題-WordCount-</h2>
|
|
|
674 </div>
|
|
|
675
|
|
|
676 <div class='slide'>
|
|
|
677 <h2>実験に利用する例題-FFT-</h2>
|
|
|
678 </div>
|
|
|
679
|
|
|
680 <div class='slide'>
|
|
|
681 <h2>実験環境</h2>
|
|
|
682 </div>
|
|
|
683
|
|
|
684 <div class='slide'>
|
|
|
685 <h2>マルチコア CPU による並列実行のベンチマーク</h2>
|
|
|
686 </div>
|
|
|
687
|
|
|
688 <div class='slide'>
|
|
|
689 <h2>DMA の prefecth に関するベンチマーク </h2>
|
|
|
690 </div>
|
|
|
691
|
|
|
692 <div class='slide'>
|
|
|
693 <h2>GPGPU のベンチマーク</h2>
|
|
|
694 </div>
|
|
|
695
|
|
|
696 <div class='slide'>
|
|
|
697 <h2>データ並列実行のベンチマーク</h2>
|
|
|
698 </div>
|
|
|
699
|
|
|
700 <div class='slide'>
|
|
|
701 <h2>GPGPU のベンチマーク</h2>
|
|
|
702 </div>
|
|
|
703
|
|
|
704 <div class='slide'>
|
|
|
705 <h2>FFT による GPGPU のベンチマーク</h2>
|
|
|
706 </div>
|
|
|
707
|
|
|
708 <div class='slide'>
|
|
|
709 <h2>BlockedRead による並列 I/O のベンチマーク</h2>
|
|
|
710 </div>
|
|
|
711
|
|
|
712 <div class='slide'>
|
|
|
713 <h2>まとめ</h2>
|
|
|
714 </div>
|
|
|
715
|
|
|
716 <div class='slide'>
|
|
|
717 <h2>今後の課題</h2>
|
|
58
|
718 </div>
|
|
|
719
|
|
|
720 </div> <!-- presentation -->
|
|
|
721 </bodypp>
|
|
|
722 </html>
|
