Mercurial > hg > Papers > 2014 > kkb-sigos
view presen/s6/index.html @ 9:9bb5fca606b3 default tip
fix fontsize
| author | Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp> |
|---|---|
| date | Tue, 13 May 2014 16:52:03 +0900 |
| parents | a5fc2e56aef7 |
| children |
line wrap: on
line source
<!DOCTYPE html> <html> <head> <meta charset='utf-8'> <title>sigos</title> <!-- Notes on CSS media types used: 1) projection -> slideshow mode (display one slide at-a-time; hide all others) 2) screen -> outline mode (display all slides-at-once on screen) 3) print -> print (and print preview) Note: toggle between projection/screen (that is, slideshow/outline) mode using t-key Questions, comments? - send them along to the mailinglist/forum online @ http://groups.google.com/group/webslideshow --> <!-- style sheet links --> <link rel="stylesheet/less" href="themes/blank/projection.css.less" media="screen,projection"> <link rel="stylesheet/less" href="themes/blank/screen.css.less" media="screen"> <link rel="stylesheet/less" href="themes/blank/print.css.less" media="print"> <link rel="stylesheet/less" href="themes/blank/blank.css.less" media="screen,projection"> <!-- Notes about less css support - all less stylesheets (*.css.less) need to get listed/loaded first (before the less.js script) - find more info about less.js online @ http://lesscss.org ***** NOTE: less.js browser script currently won’t work if you’re using Google Chrome and the path to your page starts with "file:///" due to a known Chrome issue. (In the developer/js console you will see: XMLHttpRequest cannot load file:///../s6/shared/projection.css.less. Cross origin requests are only supported for HTTP.) --> <!-- add js libs (less, jquery) --> <script src="js/less-1.1.4.min.js"></script> <script src="js/jquery-1.7.min.js"></script> <!-- S6 JS --> <script src="js/jquery.slideshow.js"></script> <script src="js/jquery.slideshow.counter.js"></script> <script src="js/jquery.slideshow.controls.js"></script> <script src="js/jquery.slideshow.footer.js"></script> <script src="js/jquery.slideshow.autoplay.js"></script> <script> $(document).ready( function() { Slideshow.init(); // Example 2: Start Off in Outline Mode // Slideshow.init( { mode: 'outline' } ); // Example 3: Use Custom Transition // Slideshow.transition = transitionScrollUp; // Slideshow.init(); // Example 4: Start Off in Autoplay Mode with Custom Transition // Slideshow.transition = transitionScrollUp; // Slideshow.init( { mode: 'autoplay' } ); } ); </script> <!-- Better Browser Banner for Microsoft Internet Explorer (IE) --> <!--[if IE]> <script src="js/jquery.microsoft.js"></script> <![endif]--> </head> <body> <div class="layout"> <div id="header"></div> <div id="footer"> <div align="right"> <img src="images/concurrency.png" width="200"> </div> </div> </div> <div class="presentation"> <!-- add slides here; example --> <div class='slide cover'> <table width="90%" height="90%" border="0" align="center"> <tr> <td><div align="center"> <h1><font color="#808db5">並列プログラミングフレームワーク Cerium の OpenCL, CUDA 対応</font></h1> </div></td> </tr> <tr> <td><div align="left"> Shohei KOKUBO <script> var date = new Date(); var year = date.getFullYear(); var month = date.getMonth(); var day = date.getDate(); var monthList = new Array("January","February","March","April","May","June", "July","August","September","October","November","December"); document.write(monthList[month]+" "+day+", "+year); </script> <hr style="color:#ffcc00;background-color:#ffcc00;text-align:left;border:none;width:300%;height:0.2em;"> </div></td> </tr> <tr> <td><div align="left" style="margin-top:-5%"> 担当教官：河野真治 </div></td> </tr> </table> </div> <div class='slide'> <h2>はじめに(1/2)</h2> <p> GPU, Cell, SpursEngine, Xeon Phi など様々なプロセッサが存在する。それぞれのプロセッサを利用するためには異なる API を用いる必要がある。それらの対応に多くの時間を取られてしまいプログラムの性能改善に集中することができない。 </p> <p> 当研究室で開発・改良が行われている並列プログラミングフレームワーク Cerium は様々なプロセッサを統合して扱えるフレームワークを目指してる。 </p> <p> Mac OS X 上で GPGPU を行うには、OpenCL または CUDA を利用する方法が考えれる。GPGPU では通常のマルチコア CPU とは<font color="red">異なる並列プログラミング</font>と<font color="red">特別なチューニング</font>が必要となる。そこで Cerium を用いてその差を吸収し、自動的なチューニングを可能にする。 </p> </div> <div class='slide'> <h2>はじめに(2/2)</h2> <p> GPU や Cell などメモリ空間が異なるプロセッサはデータ転送がオーバーヘッドとなる。このオーバーヘッドを解消するためには<font color="red">パイプライン処理</font>を行い、<font color="red">データ転送をオーバーラップ</font>する必要がある。 </p> <p> 様々なプロセッサを統合して扱えるフレームワークとして INRIA が開発している StarPU がある。StarPU は Cerium と同じタスクベースの非同期処理を行うフレームワークである。StarPU はパイプライン処理をサポートしていないので、パイプライン処理を行いたい場合は自分で実装するしかない。 </p> <p> 今回、Ceirum に OpenCL, CUDA を用いた Scheduler を実装した。 </p> <p> OpenCL, CUDA によるすべての最適化を Cerium では自動で行う。その一環として、実装した Scheduler はパイプライン処理によるデータ転送の効率化を自動で行うようになっている。 </p> </div> <div class='slide'> <h2>GPU のアーキテクチャ</h2> <table border="0" cellpadding="0" cellspacing="0"> <tbody> <tr> <td><img src="images/gpu_arch.png" style="height:350px"></td> <td> <p style="font-size:100%"> CPU と GPU はメモリ空間が異なるのでメモリの共有ができない。データにアクセスするためには空間ごとコピーする必要がある。この転送がオーバーヘッドとなる。 </p> <p style="font-size:100%"> つまり、GPGPU では<font color="red">データ転送を頻繁に行う</font>タスクは並列度が低くなる。 </p> </td> </tr> </tbody> </table> </div> <div class='slide'> <h2>並列化手法とデータ転送</h2> <p> 並列化手法は大まかにタスク並列とデータ並列に分けることができる。 </p> <p> タスク並列では複数のプロセッサで異なる演算を行い、データ並列では複数のプロセッサで同じ演算を行う。 </p> <p> タスク並列では、同じ演算を行う場合でもタスクの数だけデータ転送を行う必要がある。 </p> <p> データ並列では、すべてのデータを一度で転送する。タスク側で担当する範囲を計算し、その部分に対して処理を行う。 </p> <p> 同じ演算を行いたい場合、データ転送が少なくなるのでデータ並列の方が有効である。特に GPGPU ではデータ転送がオーバーヘッドとなるのでその差が顕著に現れる。 </p> </div> <div class='slide'> <h2>OpenCL, CUDA による GPGPU への対応</h2> <p> OpenCL では CommandQueue、CUDA では Stream という同期機構がある。それぞれに命令を投入することで GPU を制御する。 </p> <p> CommandQueue, Stream に投入された命令は発行された順序で実行される。これらを複数用いることで<font color="red">命令を並列に実行する</font>ことが可能になる。 </p> <p> OpenCL および CUDA ではタスクを kernel という単位で記述する。 </p> <p> Cerium において、CPU/Cell 上で実行する場合と GPU 上で実行する場合、それぞれ<font color="red">ほぼ同じ形式の記述</font>が可能である。 </p> </div> <div class='slide'> <h2>複数の同期機構を用いた並列実行</h2> <p> CommandQueue, Stream に投入された命令は投入された順序で実行されることが保証される。 </p> <p> 異なる CommandQueue, Stream の命令の依存関係が解消され実行可能な状態の場合、命令は並列に実行される。 </p> <img src="images/parallel_exec.png" style="height:190px"> <p> HtoD, kernel, DtoH は cuMemcpyHtoD, cuLaunchKernel, cuMemcpyDtoH という CUDA の API である。ホストからデバイスへのデータ転送、kernel の実行、デバイスからホストへのデータ転送をそれぞれ行う。OpenCL では clEnqueueReadBuffer, clEnqueueTask, clEnqueueWriteBuffer に相当する。 </p> </div> <div class='slide'> <h2>命令の投入方法</h2> <p> 注意点として CUDA では以下のように命令を投入しないと命令が並列に実行されない。 </p> <pre style="font-size:100%"> // BAD for (int i=0;i<stream_num;i++) { cuMemcpyHtoDAsync(); cuLaunchKernel(); cuMemcpyDtoHAsync(); }</pre> <pre class="code" style="font-size:100%"> // GOOD for (int i=0;i<stream_num;i++) cuMemcpyHtoDAsync(); for (int i=0;i<stream_num;i++) cuLaunchKernel(); for (int i=0;i<stream_num;i++) cuMemcpyDtoHAsync();</pre> </div> <div class='slide'> <h2>host の記述</h2> <p> 2つの input の積を取り、output に返す例題の host </p> <pre style="font-size:90%"> void multi_init(TaskManager *manager) { A = new float[length]; B = new float[length]; C = new float[length]; for(int i=0; i<length; i++) { A[i]=(float)(i+1000); B[i]=(float)(i+1)/10.f; } HTask* multiply = manager->create_task(MULTIPLY_TASK); multiply->set_cpu(CPU_TYPE); multiply->set_inData(0,(memaddr)A, sizeof(float)*length); multiply->set_inData(1,(memaddr)B, sizeof(float)*length); multiply->set_outData(0,(memaddr)C, sizeof(float)*length); multiply->iterate(length); }</pre> <p> CPU, GPU ともに完全に同じ記述で Task を生成することができる。set_cpu で実行するデバイスを選択することができる。 </div> <div class='slide'> <h2>kernel の記述</h2> <p> 2つの input の積を取り、output に返す例題の kernel </p> <pre style="font-size:90%"> // MultiCore CPU static int run(SchedTask *s) { float* A = (float*)s->get_input(0); float* B = (float*)s->get_input(1); float* C = (float*)s->get_output(0); long i = (long)s->x; C[i]=A[i]*B[i]; return 0; }</pre> <pre class="code" style="font-size:90%"> // OpenCL __kernel void multi(__global const long *params, __global const float* A, __global const float* B, __global float* C) { long id = get_global_id(0); C[id]=A[id]*B[id]; }</pre> <pre class="code" style="font-size:90%"> // CUDA extern "C" { __global__ void multi(long* params, float* A, float* B, float* C) { int id = blockIdx.x * blockDim.x + threadIdx.x; C[id]=A[id]*B[id]; } }</pre> <p> 修飾子などが多少異なるが、ほぼ同じ形式で記述できる。 </p> </div> <div class='slide'> <h2>データ並列</h2> <table border="0" cellpadding="0" cellspacing="0"> <tbody> <tr> <td><img src="images/ndrange_arch.png" style="height:350px"></td> <td> <p style="font-size:100%"> 1つの記述から ID の異なる複数の kernel を自動生成する。 </p> <p style="font-size:100%"> OpenCL では get_global_id 関数、CUDA では組み込み変数を参照することで ID を取得することができる。 </p> <p style="font-size:100%"> 取得した ID から担当する範囲を計算し、その部分に対して処理を行う。 </p> </td> </tr> </tbody> </table> </div> <div class='slide'> <h2>Cerium Task Manager(1/2)</h2> <table border="0" cellpadding="0" cellspacing="0"> <tbody> <tr> <td><img src="images/createTask.png" style="height:350px"></td> <td> <p style="font-size:100%"> TaskManager で依存関係が解消され、実行可能になった Task は ActiveTaskList に移される。 </p> <p style="font-size:100%"> ActiveTaskList に移された Task は依存関係が存在しないのでどのような順序で実行しても問題ない。 </p> <p style="font-size:100%"> Task は Scheduler に転送しやすい TaskList に変換され、Synchronized Queue である mail を通して担当する Scheduler に転送される。 </p> </td> </tr> </tbody> </table> </div> <div class='slide'> <h2>Cerium Task Manager(2/2)</h2> <table border="0" cellpadding="0" cellspacing="0"> <tbody> <tr> <td><img src="images/scheduler.png" style="height:300px"></td> <td> <p style="font-size:100%"> Scheduler ではパイプラインで Task が処理される。 </p> <p style="font-size:100%"> Task が終了すると Scheduler から TaskManager に mail を通して通知される。 </p> <p style="font-size:100%"> 通知に従って依存関係が処理され、再び TaskManager から Schdueler に Task が転送される。 </p> </td> </tr> </tbody> </table> </div> <div class='slide'> <h2>CPU におけるデータ並列のための API</h2> <table border="0" cellpadding="0" cellspacing="0"> <tbody> <tr> <td><img src="images/iterate.png" style="height:300px"></td> <td> <p style="font-size:100%"> データ並列で実行する API として「iterate」を実装した。 </p> <ul> <li style="font-size:100%">CPU におけるデータ並列</li> <ul> <li style="font-size:100%">Task を複製する。</li> <li style="font-size:100%">複製した Task に index と input/output data を割り当て、各 CPU に転送する。</li> <li style="font-size:100%">各 CPU にも ID を割り当てる。</li> <li style="font-size:100%">各 CPU に割り当てられた ID から実行する Task を選択する。</li> <li style="font-size:100%">Task 側では SchedTask を参照することで index を取得することができる。</li> </ul> </ul> </td> </tr> </tbody> </table> </div> <div class='slide'> <h2>ベンチマーク</h2> <ul> <li>測定環境</li> <ul> <li>OS：MacOS 10.9.2</li> <li>CPU：2*2.66GHz 6-Core Intel Xeon</li> <li>GPU：NVIDIA Quadro K5000 4096MB</li> <li>Memory：16GB 1333MHz DDR3</li> <li>Compiler：Apple LLVM version 5.1(clang-503.0.40)(based on LLVM 3.4svm)</li> </ul> </ul> <ul> <li>例題</li> <ul> <li>WordCount</li> <li>FFT</li> </ul> </ul> </div> <div class='slide'> <h2>WordCount(1/2)</h2> <table> <tbody> <tr> <td> <!-- benchmark --> <table border="2" style="font-size:90%"> <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;">0.73s</td> </tr> <tr> <td style="text-align: center;">2 CPU</td> <td style="text-align: right;">0.38s</td> </tr> <tr> <td style="text-align: center;">4 CPU</td> <td style="text-align: right;">0.21s</td> </tr> <tr> <td style="text-align: center;">8 CPU</td> <td style="text-align: right;">0.12s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">OpenCL(no pipeline)</td> <td style="text-align: right; " bgcolor="#ffffcc">48.32s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">OpenCL(pipeline)</td> <td style="text-align: right;" bgcolor="#ffffcc">46.74s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">CUDA(no pipeline)</td> <td style="text-align: right; " bgcolor="#ffffcc">55.71s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">CUDA(pipeline)</td> <td style="text-align: right;" bgcolor="#ffffcc">10.26s</td> </tr> </tbody> </table> </td> <td> <p style="font-size:100%"> 100MB のテキストファイルに対して<font color="red">タスク並列</font>で WordCount を行なった。 </p> <p style="font-size:100%"> パイプライン処理を行うことで CUDA では<font color="red">5.4倍</font>の性能向上が見られた。一方、OpenCL では性能向上が見られなかった。 </p> </tr> </tbody> </table> </div> <div class='slide'> <h2>WordCount(2/2)</h2> <table> <tbody> <tr> <td> <!-- benchmark --> <table border="2" style="font-size:90%"> <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;">0.73s</td> </tr> <tr> <td style="text-align: center;">2 CPU</td> <td style="text-align: right;">0.38s</td> </tr> <tr> <td style="text-align: center;">4 CPU</td> <td style="text-align: right;">0.21s</td> </tr> <tr> <td style="text-align: center;">8 CPU</td> <td style="text-align: right;">0.12s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">Data Parallel by OpenCL</td> <td style="text-align: right;" bgcolor="#ffffcc">0.38s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">Data Parallel by CUDA</td> <td style="text-align: right;" bgcolor="#ffffcc">0.71s</td> </tr> </tbody> </table> </td> <td> <p style="font-size:100%"> 100MB のテキストファイルに対して<font color="red">データ並列</font>で WordCount を行なった。 </p> <p style="font-size:100%"> データ並列による実行は 1CPU に対して OpenCL では<font color="red">1.9倍</font>、CUDA では<font color="red">ほぼ同等</font>という結果になった。 </p> </tr> </tbody> </table> </div> <div class='slide'> <h2>FFT</h2> <table> <tbody> <tr> <td> <!-- benchmark --> <table border="2" style="font-size:90%"> <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;">0.48s</td> </tr> <tr> <td style="text-align: center;">2 CPU</td> <td style="text-align: right;">0.26s</td> </tr> <tr> <td style="text-align: center;">4 CPU</td> <td style="text-align: right;">0.17s</td> </tr> <tr> <td style="text-align: center;">8 CPU</td> <td style="text-align: right;">0.11s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">OpenCL</td> <td style="text-align: right;" bgcolor="#ffffcc">0.09s</td> </tr> <tr> <td style="text-align: center;" bgcolor="#ffffcc">CUDA</td> <td style="text-align: right;" bgcolor="#ffffcc">0.21s</td> </tr> </tbody> </table> </td> <td> <p style="font-size:100%"> 512*512 の画像ファイルに対してフーリエ変換を行い、その後、High Pass Filter をかける例題である。 </p> <p style="font-size:100%"> 1CPU に対して OpenCL では<font color="red">5.3倍</font>、CUDA では<font color="red">2.2倍</font>の性能向上が見られた。 </p> </tr> </tbody> </table> </div> <div class='slide'> <h2>考察</h2> <p> パイプライン処理を行うことで CUDA では性能向上が確認できた。しかし、OpenCL では性能向上が見られなかった。OpenCL と CUDA をそれぞれ用いた Scheduler はほぼ同等の実装である。このことから OpenCL では CUDA とは異なるチューニングが必要になると考えられる。 </p> <p> データ並列による実行はタスク並列に比べ、どちらも性能向上が確認できた。しかし、WordCount, FFT の両方で CUDA は OpenCL に劣る結果となっている。CUDA によるデータ並列実行の機構を見直す必要がある。 </p> <p> GPU のみで演算を行う場合でも1つの Task 終了時にメインメモリにデータ転送を行なっている。これは不必要なデータ転送である。GPU 上に処理されたデータがあるので、それを利用するようにデータ転送の機構を改善する必要がある。 </div> <div class='slide'> <h2>まとめ</h2> <ul> <li>Cerium を OpenCL および CUDA に対応させた。</li> <li>WordCount, FFT を例題として実装した Scheduler の評価を行なった。</li> <li>OpenCL と CUDA で異なるチューニングが必要なことを確認した。</li> </ul> <h2>今後の課題</h2> <ul> <li>OpenCL, CUDA を用いた Scheduler の同期機構の見直し</li> <li>データ転送機構の見直し</li> <li>kernel の自動変換</li> </ul> </div> </div> <!-- presentation --> </body> </html>