Mercurial > hg > Papers > 2013 > yuhi-thesis
view presen/index.html @ 8:954053cacd0e default tip
add resume
author | Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp> |
---|---|
date | Thu, 21 Feb 2013 19:29:57 +0900 |
parents | 84766a598c3a |
children |
line wrap: on
line source
<!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>