Mercurial > hg > Papers > 2013 > yuhi-prosym

comparison presen/index.html @ 29:6b60c1277599 default tip

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
fix
author Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
date Sun, 12 Jan 2014 09:18:32 +0900
parents d6d545336b5c
children
comparison
equal deleted inserted replaced
28:d6d545336b5c 29:6b60c1277599
106 106
107 <slide> 107 <slide>
108 <hgroup> 108 <hgroup>
109 <h3>Kernelの記述</h3> 109 <h3>Kernelの記述</h3>
110 </hgroup> 110 </hgroup>
111 <article>
112 2つのinputDataの積を取り、outputDataに返す例題Multiply 111 2つのinputDataの積を取り、outputDataに返す例題Multiply
113 <table border="0" cellpadding="0" cellspacing="0"> 112 <table border="0" cellpadding="0" cellspacing="0">
114 <tbody> 113 <tbody>
115 <tr> 114 <tr>
116 <td> 115 <td>
117 <img src='images/kernel_description.png' style="width:650px"> 116 <img src='images/kernel_description.png' style="height:570px">
118 </td> 117 </td>
119 <td style="font-size:18pt;color:black"> 118 <td style="font-size:18pt;color:black">
120 <p> 119 <p>
121 Taskの処理自体はC/C++の形式なので、同じ記述が可能。 120 Taskの処理自体はC/C++の形式なので、同じ記述が可能。
122 </p> 121 </p>
123 <p> 122 <p>
124 引数の受け取り方が違う。 123 引数の受け取り方が違う。
125 <ul> 124 <ul>
126 <li>CPU</li> 125 <li>CPU
127 <dd>rbuf/wbufとしてDataをまとめて受け取っている</dd> 126 <dd>rbuf/wbufとしてDataをまとめて受け取っている</dd></li>
128 <li>GPU</li> 127 <li>GPU
129 <dd>1つ1つ個別の変数として受け取っている</dd> 128 <dd>1つ1つ個別の変数として受け取っている</dd>
130 <dd>それぞれの変数にOpenCL独自の修飾子が必要</dd> 129 <dd>それぞれの変数にOpenCL独自の修飾子が必要</dd></li>
131 </ul> 130 </ul>
132 変数で受けるなどしてこの差異を吸収すれば、同じ記述が可能。 131 変数で受けるなどしてこの差異を吸収すれば、同じ記述が可能。
133 </p> 132 </p>
134 133
135 </td> 134 </td>
136 </tr> 135 </tr>
137 </tbody> 136 </tbody>
138 </table> 137 </table>
139 <br>
140 </article>
141 </slide> 138 </slide>
142 139
143 140
144 <slide> 141 <slide>
145 <hgroup> 142 <hgroup>
204 </hgroup> 201 </hgroup>
205 <article> 202 <article>
206 <table border="0" cellpadding="0" cellspacing="0"> 203 <table border="0" cellpadding="0" cellspacing="0">
207 <tbody> 204 <tbody>
208 <tr> 205 <tr>
209 <td><img src='images/gpu_data_parallel.png' style="height:400px"></td> 206 <td><img src='images/gpu_data_parallel.png' style="height:450px"></td>
210 <td style="font-size:18pt;color:black"> 207 <td style="font-size:18pt;color:black">
211 <p> 208 <p>
212 タスク並列だと、タスクごとにinput data と output data を転送しなければならない 209 タスク並列だと、タスクごとにinput data と output data を転送しなければならない
213 </p> 210 </p>
214 <p> 211 <p>
232 </hgroup> 229 </hgroup>
233 <article> 230 <article>
234 <table border="0" cellpadding="0" cellspacing="0"> 231 <table border="0" cellpadding="0" cellspacing="0">
235 <tbody> 232 <tbody>
236 <tr> 233 <tr>
237 <td><img src='images/ndrange_arch.png' style="height:450px"></td> 234 <td><img src='images/ndrange_arch.png' style="height:500px"></td>
238 <td style="font-size:18pt;color:black"> 235 <td style="font-size:18pt;color:black">
239 <p> 236 <p>
240 データを2、3次元に分割し、分割した部分に対して 237 データを2、3次元に分割し、分割した部分に対して
241 同一のTaskを割り当て、並列に処理を行う並列化手法 238 同一のTaskを割り当て、並列に処理を行う並列化手法
242 </p> 239 </p>
296 </hgroup> 293 </hgroup>
297 <article> 294 <article>
298 <p> 295 <p>
299 input data を二倍してoutput data に返す例題 296 input data を二倍してoutput data に返す例題
300 </p> 297 </p>
301 <pre class="prettyprint" data-lang="main.cc(Iterate Task 生成)"> 298 <img src="images/iterateTaskGen.png" height="250"></img>
302 HTaskPtr twice = manager->create_task(Twice);
303 twice->set_cpu(GPU);
304 twice->set_inData(0,(memaddr)input, sizeof(float)*length);
305 twice->set_outData(0,(memaddr)output, sizeof(float)*length);
306 // paramに0~length-1(index)をsetしたtaskをlength個spawnする
307 twice->iterate(length);
308 </pre>
309 <p> 299 <p>
310 iterate(length)とすることで、TaskManagerがデータ並列用にTaskを生成 300 iterate(length)とすることで、TaskManagerがデータ並列用にTaskを生成
311 </p> 301 </p>
312 <p> 302 <p>
313 この例だとlength個のTaskが生成され、各Taskに0~length-1までのID(index)を割り振る 303 この例だとlength個のTaskが生成され、各Taskに0~length-1までのID(index)を割り振る
321 <slide> 311 <slide>
322 <hgroup> 312 <hgroup>
323 <h3>iterate(Kernel)</h3> 313 <h3>iterate(Kernel)</h3>
324 </hgroup> 314 </hgroup>
325 <article> 315 <article>
326 <pre class="prettyprint" data-lang="twice.cc(MultiCore)"> 316 <img src="images/source/iterate_kernel.png" height="300"></img>
327
328 long i = (long)scheduler->x; // (long)scheduler->get_param(0);
329 output[i]=input[i]*2;
330 </pre>
331 <pre class="prettyprint" data-lang="twice.cl(GPU)">
332 long i = get_global_id(0);
333 output[i]=input[i]*2;
334 </pre>
335 <p> 317 <p>
336 MultiCoreではschedulerの持つメンバ変数x、OpenCLはget_global_idというAPIを用いて、 318 MultiCoreではschedulerの持つメンバ変数x、OpenCLはget_global_idというAPIを用いて、
337 自分に割り振られたid(index)を取得する。その後、そのindexに対して処理を行う。 319 自分に割り振られたid(index)を取得する。その後、そのindexに対して処理を行う。
338 </p> 320 </p>
339 <p> 321 <p>
345 <slide> 327 <slide>
346 <hgroup> 328 <hgroup>
347 <h3>iterateの実装</h3> 329 <h3>iterateの実装</h3>
348 </hgroup> 330 </hgroup>
349 <article> 331 <article>
350 <pre class="prettyprint" data-lang="HTask.cc"> 332 <img src="images/source/iterate.png" height="280"></img>
351 void 333 <p>渡されたlengthはwork item の要素数となる。</p>
352 HTask::iterate(long x) {
353 tl->dim=1;
354 tl->x=x;
355 // 1次元なのでy軸、z軸の要素数は1
356 tl->y=1;
357 tl->z=1;
358 mimpl->spawn_task(this);
359 }
360 </pre>
361 <p>渡されたlengthはwork item のx座標の要素数となる。</p>
362 <p>taskにデータ分割のためのdata(dimension, や workitem size)をsetする。</p> 334 <p>taskにデータ分割のためのdata(dimension, や workitem size)をsetする。</p>
363 <p> 335 <p>
364 Schedulerはtaskが持ってるdataを元にtaskを複数生成し、idを割り当てる。 336 Schedulerはtaskが持ってるdataを元にtaskを複数生成し、idを割り当てる。
365 </p> 337 </p>
366 </article> 338 </article>
369 <slide> 341 <slide>
370 <hgroup> 342 <hgroup>
371 <h3>iterateの実装(多次元)</h3> 343 <h3>iterateの実装(多次元)</h3>
372 </hgroup> 344 </hgroup>
373 <article> 345 <article>
374 <pre class="prettyprint" data-lang="HTask.cc"> 346 <img src="images/source/iterate_multidim.png" height="300"></img>
375 void
376 HTask::iterate(long x, long y) {
377 tl->dim=2;
378 tl->x=x;
379 tl->y=y;
380 tl->z=1;
381 mimpl->spawn_task(this);
382 }
383 </pre>
384 <p> 347 <p>
385 引数を複数渡せば多次元のデータ分割ができる。3次元までサポートする。 348 引数を複数渡せば多次元のデータ分割ができる。3次元までサポートする。
386 </p> 349 </p>
387 </article> 350 </article>
388 </slide> 351 </slide>
393 </hgroup> 356 </hgroup>
394 <br> 357 <br>
395 <p> 358 <p>
396 GpuSchedulerがOpenCLのAPIを呼び出し、GPUの制御を行う</p> 359 GpuSchedulerがOpenCLのAPIを呼び出し、GPUの制御を行う</p>
397 <p> 360 <p>
398 TaskManagerからTaskを受け取り、Command Queueにenqueueする 361 TaskManagerから受け取ったTaskは必要なパラメータを各種持っている
399 </p> 362 </p>
400 <article> 363 <article>
401 <pre class="prettyprint" data-lang="GpuScheduler.cc"> 364 <img src="images/source/GpuScheduler.png" height="120"></img>
402
403 clEnqueueNDRangeKernel(command_queue, kernel[cur], task->dim, NULL,
404 &task->x, &task->y, &task->z, NULL, NULL);
405 </pre>
406 <table border="2" style="font-size:18pt;"> 365 <table border="2" style="font-size:18pt;">
407 <tbody> 366 <tbody>
408 <h3 class="yellow">Taskの持つメンバ変数</h3> 367 <h3 class="yellow">Taskの持つメンバ変数</h3>
409 <tr> 368 <tr>
410 <td>dim</td> 369 <td>dim</td>
534 </tr> 493 </tr>
535 </tbody> 494 </tbody>
536 </table> 495 </table>
537 <br> 496 <br>
538 <p> 497 <p>
539 CPUとGPUの同時実行は、実行するTaskに対してset_cpu(SPE_ANY)とすれば良い 498 CPUとGPUの同時実行は、実行するTaskに対してset_cpu(ANY_ANY)とすれば良い
540 </p> 499 </p>
541 <p> 500 <p>
542 TaskはCPUとGPU、交互に割り振られる 501 TaskはCPUとGPU、交互に割り振られる
543 </p> 502 </p>
544 </article> 503 </article>
547 <slide> 506 <slide>
548 <hgroup> 507 <hgroup>
549 <h3>ベンチマーク</h3> 508 <h3>ベンチマーク</h3>
550 </hgroup> 509 </hgroup>
551 <article> 510 <article>
552 <table > 511 <img src="images/bench_mark_hetero.png" height="350"></img>
553 <tbody> 512 <h3 class="yellow">結果</h3>
554 <tr> 513 <p>
555 <td> 514 現段階ではSchedulingを行ってないため、GPUやCPUを単体で動かした時よりも遅くなる。
556 <img src="images/bench_mark_hetero.png" height="400"></img> 515 </p>
557 </td> 516 <p>
558 <td> 517 CPUとGPUは実行速度に差が出る場合がある。
559 <h3 class="yellow">結果</h3> 518 Taskの計算内容によってはどちらかのアーキテクチャに任せた方が良い。
560 <font size="5"> 519 それを確認するベンチマークを動かす。
561 <p> 520 </p>
562 現段階ではSchedulingを行ってないため、GPU単体で動かした時よりも遅くなる。 521 </font>
563 </p> 522 </article>
564 <p> 523 </slide>
565 CPUとGPUは実行速度に差が出る場合がある。 524
566 Taskの計算内容によってはどちらかのアーキテクチャに任せた方が良い。 525
567 それを確認するベンチマークを動かす。 526 <slide>
568 </p> 527 <hgroup>
569 <p> 528 <h3>ベンチマーク</h3>
570 529 </hgroup>
571 </p> 530 <article>
572 </font> 531 <img src="images/bench_mark_each_task.png" height="350"></img>
573 </td>
574 </tr>
575 </tbody>
576 </table>
577 </article>
578 </slide>
579
580
581 <slide>
582 <hgroup>
583 <h3>ベンチマーク</h3>
584 </hgroup>
585 <article>
586 <img src="images/bench_mark_each_task.png" height="350"></img>
587 <p> 532 <p>
588 FFTはSpinFactやButterfly演算等、様々なTaskで構成されている。 533 FFTはSpinFactやButterfly演算等、様々なTaskで構成されている。
589 それぞれのTaskについて、実行時間を計測した。 534 それぞれのTaskについて、実行時間を計測した。
590 </p> 535 </p>
591 <p> 536 <p>
592 SpinFactのTaskに関しては、CPUの方が実行速度が早い。 537 大体のTaskはGPUの方が速いが、SpinFactのTaskに関してはCPUの方が実行速度が速い。
593 </p> 538 </p>
594 </article> 539 </article>
595 </slide> 540 </slide>
596 541
597 <slide> 542 <slide>
603 <tbody> 548 <tbody>
604 <tr> 549 <tr>
605 <td> 550 <td>
606 <img src="images/decide_weight.png" height="150"></img> 551 <img src="images/decide_weight.png" height="150"></img>
607 </td> 552 </td>
608 <td> 553 <td style="font-size:18pt;color:black">
609 並列実行するTaskをCPUとGPUで事前に一度実行し、実行時間を測定する。 554 並列実行するTaskをCPUとGPUで事前に一度実行し、実行時間を測定する。
610 それぞれの実行時間の割合で重みをつける。 555 それぞれの実行時間の割合で重みをつける。
611 </td> 556 </td>
612 </tr> 557 </tr>
613 <tr> 558 <tr>
614 <td> 559 <td>
615 <img src="images/select_arch.png" height="200"></img> 560 <img src="images/select_arch.png" height="200"></img>
616 </td> 561 </td>
617 <td> 562 <td style="font-size:18pt;color:black">
618 <p> 563 <p>
619 それぞれの重みからCPU実行とGPU実行のどちらに適しているか判断する 564 それぞれの重みからCPU実行とGPU実行のどちらに適しているか判断する
620 </p> 565 </p>
621 </td> 566 </td>
622 </tr> 567 </tr>
638 <tbody> 583 <tbody>
639 <tr> 584 <tr>
640 <td> 585 <td>
641 <img src="images/decide_weight2.png" height="150"></img> 586 <img src="images/decide_weight2.png" height="150"></img>
642 </td> 587 </td>
643 <td> 588 <td style="font-size:18pt;color:black">
644 全てのTaskがCPUの二倍、GPUの方が実行速度が早い場合 589 全てのTaskがCPUの二倍、GPUの方が実行速度が早い場合
645 </td> 590 </td>
646 </tr> 591 </tr>
647 <tr> 592 <tr>
648 <td> 593 <td>
649 <img src="images/select_arch2.png" height="180"></img> 594 <img src="images/select_arch2.png" height="180"></img>
650 </td> 595 </td>
651 <td> 596 <td style="font-size:18pt;color:black">
652 <p> 597 <p>
653 それぞれのTaskを得意とするアーキテクチャに全て割り振るのではなく、 598 それぞれのTaskを得意とするアーキテクチャに全て割り振るのではなく、
654 RunTimeが最小になるように割り振る 599 RunTimeが最小になるように割り振る
655 </p> 600 </p>
656 </td> 601 </td>
672 <li>CPUとGPUでのTaskの同時実行に対応</li> 617 <li>CPUとGPUでのTaskの同時実行に対応</li>
673 <li>同時実行時のTaskのScheduling手法の提案</li> 618 <li>同時実行時のTaskのScheduling手法の提案</li>
674 </ul> 619 </ul>
675 <h3 class="yellow">今後の課題</h3> 620 <h3 class="yellow">今後の課題</h3>
676 <ul> 621 <ul>
677 <li>提案したSchedulingの手法を実装・ベンチマーク</li> 622 <li>提案したSchedulingの手法を実装/ベンチマーク</li>
678 <li>ベンチマークに使用する例題の追加</li> 623 <li>ベンチマークに使用する例題の追加</li>
624 <li>GPUのSchedulerにパイプライン機構の導入</li>
679 </ul> 625 </ul>
680 </article> 626 </article>
681 </slide> 627 </slide>
682 628
683 <!-- 629 <!--