Mercurial > hg > Papers > 2014 > masakoha-thesis > final
diff slide/index.html @ 13:ff5c2b4ca18a
add slide
| author | Masataka Kohagura <e085726@ie.u-ryukyu.ac.jp> |
|---|---|
| date | Tue, 28 Jan 2014 00:48:21 +0900 |
| parents | |
| children | 0e7c972b5ca1 |
line wrap: on
line diff
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/slide/index.html Tue Jan 28 00:48:21 2014 +0900 @@ -0,0 +1,388 @@ +<!DOCTYPE html> + +<html> + <head> + <title>Presentation</title> + + <meta charset='utf-8'> + <script + src='./slides.js'></script> + </head> + + <style> + /* Your individual styles here, or just use inline styles if that’s + what you want. */ + + </style> + + <body style='display: none'> + <section class='slides layout-regular template-default'> + + <!-- + Your slides (<article>s) go here. Delete or comment out the + slides below. + --> + <article > + <h1>Cerium TaskManagerのGPGPUのサポート</h1> + <h3 class="title">Yuhi TOMARI 21 Feb 2013</h3> + <div align="right">担当教官 : 河野 真治</div> + </article> + <article > + <h3>研究概要</h3> + <p> + 当研究室ではCellおよびLinux、 + Mac OSX上で動く並列プログラミングフレームワーク、 + Ceriumの開発・改良を行っている + </p> + <img src='images/cerium_resource.png' style="height:300px"> + <p> + <font color="#ff3300">本研究では新たにGPU上での並列実行に対応</font> + </p> + <p> + これによりヘテロジニアス(異種混合)環境下<br>でのプログラミングをサポートする + </p> + + <!-- + <p> + CeriumにOpenCLを組み込む事でGpu上での並列実行にも対応すると、 + CPU・GPUそれぞれにtaskを割り振ることが可能となる。</p> + <p> + これにより、ヘテロジニアス(異種混合)な環境で + それぞれのアーキテクチャに適したプログラミングをサポートする。 + </p> + --> + + </article> + <article> + <h3>OpenCLによるGPGPUへの対応</h3> + <p> + OpenCLは + ヘテロジニアス環境での並列計算を支援するフレームワーク + </p> + <table border="0" cellpadding="0" cellspacing="0"> + <tbody> + <tr> + <td><img src='images/opencl.png' style="height:120px"></td> + <td> + <p> + OpenCLのtaskはkernelと呼ばれ、OpenCL Cという + Cに似た独自の言語で記述される + </p> + </td> + </tr> + </tbody> + </table> + <p> + <h3 class='yellow'>CPU/Cell上で実行する場合とGPU上で実行する場合、<br> + それぞれほぼ同じ形式でkernelを記述できる</h3> + </p> + </article> + <article> + <h3>Cerium Task Manager</h3> + <table border="0" cellpadding="0" cellspacing="0"> + <tbody> + <tr> + <td><img src='images/createtask.png' style="height:350px"></td> + <td> + <ol> + <li>Taskを生成</li> + <li>依存関係のチェック</li> + <li>Schedulerに転送</li> + <li>並列実行</li> + </ol> + </td> + </tr> + </tbody> + </table> + <p> + CpuThreads、Schedulerに対応させる形でGpuThreadsとGpuSchedulerを作成した + </p> + </article> + <article class="nobackground"> + <h3>GPUのアーキテクチャ</h3> + <table border="0" cellpadding="0" cellspacing="0"> + <tbody> + <tr> + <td><img src='images/gpu_arch.png' style="height:400px"></td> + <td> + <p> + GPU上のローカルメモリはCPUとメモリ空間が異なるのでメモリの共有ができない。 + </p> + </td> + </tr> + </tbody> + </table> + <p> + メモリのコピー処理やTaskを順にCommand Queueに + enqueueする。enqueueされた処理は並列実行される + </article> + <article> + <h3>GPU Task実行の流れ</h3> + <br> + <h3 class="yellow">kernel fileの記述</h3> +<pre>__kernel void // kernel.cl(kernel file) +twice(__global int *input_data,__global int *output_data) { + long count = (long)data_count[0]; + for (int i = 0; i< count; i++) { + output_data[i] = input_data[i] * 2;; + } +} +</pre> + </article> + <article> + <h3>GPU Task 実行の流れ</h3> + <br> + <h3 class="yellow">kernelをTaskとしてCeriumに登録</h3> + <pre>void +task_init(void) { // task_init.cc + GpuSchedRegister(Twice, "./twice.cl", "twice"); +}</pre> + <table border="2" style="font-size:18pt;"> + <tbody> + <tr> + <td> 第1引数<br>Twice</td> + <td >Taskのid。enumで定義されている<br>TaskManagerはこの値でTaskを識別する</td> + </tr> + <tr> + <td> 第2引数<br>"./twice.cl"</td> + <td>OpenCLが処理するkernelが記述されているファイルのパス</td> + </tr> + <tr> + <td> 第3引数<br>"twice"</td> + <td >関数の指定。kernel file内にある、実行する関数名を指定<br> + Taskにあたる部分</td> + </tr> + </tbody> + </table> + </article> + <article class="nobackground"> + <h3>GPU Task 実行の流れ</h3><br> + <h3 class="yellow">GPU Task生成</h3> +<pre>// main.cc +HTaskPtr twice = manager->create_task(Twice); +twice->set_inData(0, indata, sizeof (int)*length); +// twice->set_inData(1, indata2, sizeof (int)*length); +twice->set_outData(0, outdata, sizeof (int)*length); +twice->set_cpu(GPU); +twice->spawn(); + </pre> + </article> + + + <article> + <h3>Cerium OpenCL API比較</h3> + <img src='images/api.png' style="height:500px"> + </article> + <article class="nobackground"> + <h3>ベンチマーク</h3> + <table > + <tbody> + <tr> + <td> <!-- benchmark --> + <table border="2" style="font-size:18pt;"> + <tbody> + <tr> + <td bgcolor="#8091B5"></td> + <td style="text-align: center;">Time</td> + </tr> + <tr> + <td style="text-align: center;">1 CPU</td> + <td style="text-align: right;"> 796 ms</td> + </tr> + <tr> + <td style="text-align: center;">2 CPU</td> + <td style="text-align: right;"> 439 ms</td> + </tr> + <tr> + <td style="text-align: center;">6 CPU</td> + <td style="text-align: right;"> 153 ms</td> + </tr> + <tr> + <td style="text-align: center;">12 CPU</td> + <td style="text-align: right;"> 96 ms</td> + </tr> + <tr> + <td style="text-align: center;">24 CPU</td> + <td style="text-align: right;"> 89 ms</td> + </tr> + <tr> + <td style="text-align: center;" bgcolor="#ffffcc">GPU(改良前)</td> + <td style="text-align: right; " bgcolor="#ffffcc"> 330752 ms</td> + </tr> + <tr> + <td style="text-align: center;" bgcolor="#ffffcc">GPU(改良後)</td> + <td style="text-align: right;" bgcolor="#ffffcc"> 5306 ms</td> + </tr> + </tbody> + </table> + </td> <!-- /benchmark --> + <td> <!-- system env --> + <h3 class="yellow">10万入力によるBitonic Sort</h3> + <font size="5"> + <p>実験環境</p> + OS : MacOS 10.8.2<br> + CPU : 2*2.66GHz 6-CoreIntel Xeon<br> + Memory : 16GB<br> + Compiler : Apple clang version 4.1 <br> (based on LLVM 3.1svn)<br> + GPU : AMD ATI Radeon HD 5870 1024MB<br> + </font> + <h3 class="yellow">結果</h3> + 1coreのCPUよりも10倍遅い + </td> <!--system env --> + </tr> + </tbody> + </table> + <p> + 充分な性能が出なかったため、一度に送信する + data のサイズを増やす改善を行ったところ、 + <font color="red">約60倍</font>程実行速度が向上した + </p> + </article> + <article> + <h3>考察</h3> + <p> + 性能向上は見られたが、CPUと比べると未だ差が開いている + GPU向けに適切なチューニングが今後の課題となる + </p><br> + <h3 class="yellow">改善案</h3> + <ul> + <li>データ並列によるkernelの実行</li> + <li>同期機構の見直し</li> + </ul> + </article> + <article> + <h3>データ並列</h3> + <p> + データを2、3次元に分割し、分割した部分に対して並列処理する並列化手法。 + </p> + <p> + OpenCL ではin/outするデータ郡をWork Itemと呼ぶ。 + </p> + <table border="0" cellpadding="0" cellspacing="0"> + <tbody> + <tr> + <td><img src='images/ndrange_arch.png' style="height:350px"></td> + <td>各Work Item のサイズを指定するとOpenCLがデータ並列で実行する。</td> + </tr> + </tbody> + </table> + </article> + <article> + <h3>同期機構</h3> + <p> + GpuSchedulerはCommand Queueの内部でパイプライン的に実行を行っている。 + パイプラインを構成するには処理にwaitをかける必要がある。現在はclWaitForEvent APIを使用 + </p> + <table border="0" cellpadding="0" cellspacing="0"> + <tbody> + <tr> + <td bgcolor="#8091B5"><font color="white">API</font></td> + <td>機能</td> + </tr> + <tr> + <td bgcolor="#8091B5"><font color="white">clFlush()</font></td> + <td>Command Queueに投入したTask全てをDeviceで実行する</td> + </tr> + <tr> + <td bgcolor="#8091B5"><font color="white">clWaitForEvent()</font></td> + <td>特定の処理の終了を待つ</td> + </tr> + </tbody> + </table> + <p>clFlushは実行は保証するが、<font color="red">終了は保証しない</font>仕様になっている</p> + </article> + <article> + <h3>新しい同期</h3> + <table border="0" cellpadding="0" cellspacing="0"> + <tbody> + <tr> + <td bgcolor="#8091B5"><font color="white">FrameWork</font></td> + <td>Dependency</td> + </tr> + <tr> + <td bgcolor="#8091B5"><font color="white">Cerium</font></td> + <td>Task Dependency</td> + </tr> + <tr> + <td bgcolor="#8091B5"><font color="white">OpenCL</font></td> + <td>Data Dependency</td> + </tr> + </tbody> + </table> + <p> + Task Dependency:Schedulerで依存関係が決定<br> + Data Dependency:GPUに読み込まれた時に決定 + </p> + <p> + GPGPUはなるべくGPU内部で処理を行う方が高速なため、性能向上が見込める + </p> + </article> + <article> + <h3>まとめ</h3> + <ul> + <li>Cerium Task ManagerをGPGPUに対応</li> + <li>同期機構の実装</li> + <li>マルチコア実行とGPU実行のベンチマーク</li> + </ul> + <h3 class="yellow">今後の課題</h3> + <ul> + <li>データ並列による実行のサポート</li> + <li>同期機構の見直し</li> + </ul> + </article> + <article> + <h3>ベンチマーク</h3> + <table > + <tbody> + <tr> + <td> <!-- benchmark --> + <table border="2" style="font-size:18pt;"> + <tbody> + <tr> + <td bgcolor="#8091B5"></td> + <td style="text-align: center;">Time</td> + </tr> + <tr> + <td style="text-align: center;">1 CPU</td> + <td style="text-align: right;"> 67 ms</td> + </tr> + <tr> + <td style="text-align: center;">2 CPU</td> + <td style="text-align: right;"> 34 ms</td> + </tr> + <tr> + <td style="text-align: center;">6 CPU</td> + <td style="text-align: right;"> 12 ms</td> + </tr> + <tr> + <td style="text-align: center;">12 CPU</td> + <td style="text-align: right;"> 9 ms</td> + </tr> + <tr> + <td style="text-align: center;">24 CPU</td> + <td style="text-align: right;"> 6 ms</td> + </tr> + <tr> + <td style="text-align: center;" bgcolor="#ffffcc">GPU</td> + <td style="text-align: right;" bgcolor="#ffffcc"> 10201 ms</td> + </tr> + </tbody> + </table> + </td> <!-- /benchmark --> + <td> <!-- system env --> + <h3 class="yellow">word count</h3> + <font size="5"> + <p>10MBのテキストファイルを分割<br> + 各Taskがcountしていく</p> + <p>スペースと改行区切りでword countしていく</p> + </font> + <h3 class="yellow">結果</h3> + CPUの方が150倍早い + </td> <!--system env --> + </tr> + </tbody> + </table> + </article> + <!--- <img src='images/flow_chart.jpg' width="300" height="500"> --> + </body> +</html>