Mercurial > hg > Papers > 2014 > masakoha-thesis > final
changeset 15:9b071b32e3de
add some files and write slide
| author | Masataka Kohagura <e085726@ie.u-ryukyu.ac.jp> |
|---|---|
| date | Wed, 29 Jan 2014 16:12:51 +0900 |
| parents | 0e7c972b5ca1 |
| children | 56cd6800acb5 |
| files | paper/chapter2.tex paper/fig/createTask1.bb paper/fig/createTask1.pdf paper/fig/fig.graffle slide/OUTLINE slide/example.html slide/images/cerium.png slide/images/createTask.graffle slide/images/createTask.png slide/index.html |
| diffstat | 10 files changed, 1990 insertions(+), 281 deletions(-) [+] |
line wrap: on
line diff
--- a/paper/chapter2.tex Tue Jan 28 18:17:39 2014 +0900 +++ b/paper/chapter2.tex Wed Jan 29 16:12:51 2014 +0900 @@ -49,7 +49,7 @@ \hline set\_cpu & Task を実行するデバイスの設定 \\ \hline - spawn& 生成した Task を ActiveTaskList に登録 \\ + spawn & 生成した Task を TaskList に set \\ \hline \end{tabular} \end{center}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/paper/fig/createTask1.bb Wed Jan 29 16:12:51 2014 +0900 @@ -0,0 +1,5 @@ +%%Title: ./fig/createTask1.pdf +%%Creator: extractbb 20120420 +%%BoundingBox: 0 0 431 380 +%%CreationDate: Sat Jan 25 13:46:54 2014 +
--- a/paper/fig/fig.graffle Tue Jan 28 18:17:39 2014 +0900 +++ b/paper/fig/fig.graffle Wed Jan 29 16:12:51 2014 +0900 @@ -46,7 +46,7 @@ <key>Creator</key> <string>MasaKoha</string> <key>DisplayScale</key> - <string>1 0/72 in = 1 0/72 in</string> + <string>1 0/72 in = 1.0000 in</string> <key>GraphDocumentVersion</key> <integer>8</integer> <key>GraphicsList</key> @@ -94,8 +94,8 @@ {\colortbl;\red255\green255\blue255;} \pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural -\f0\fs24 \cf0 1:createtask()\ -2:spawn()}</string> +\f0\fs24 \cf0 createtask()\ +spawn()}</string> <key>VerticalPad</key> <integer>0</integer> </dict> @@ -178,8 +178,8 @@ <integer>51</integer> <key>Points</key> <array> - <string>{289.49998931525181, 249.33334959366042}</string> - <string>{333.50003704289043, 249.62094569373764}</string> + <string>{289.49999909181315, 249.09607114673688}</string> + <string>{333.5000030945593, 249.17895606666326}</string> </array> <key>Style</key> <dict> @@ -211,8 +211,8 @@ <integer>50</integer> <key>Points</key> <array> - <string>{130.40079054872515, 210.16609773774039}</string> - <string>{194.74146722634595, 232.82665881410423}</string> + <string>{129.88006358058934, 210.16780584070199}</string> + <string>{193.49136532333424, 232.83095557438625}</string> </array> <key>Style</key> <dict> @@ -244,8 +244,8 @@ <integer>49</integer> <key>Points</key> <array> - <string>{189.65433012237426, 156.16959324794502}</string> - <string>{129.57684107816425, 177.83118429683719}</string> + <string>{189.59775698691453, 156.16940234575466}</string> + <string>{129.44447355269742, 177.83073988701349}</string> </array> <key>Style</key> <dict> @@ -343,8 +343,8 @@ <integer>46</integer> <key>Points</key> <array> - <string>{285.49999984893446, 72.041262911437741}</string> - <string>{333.50000060542584, 72.080098593518002}</string> + <string>{285.50000000038938, 72.010950396119398}</string> + <string>{333.50000003776938, 72.021256651298529}</string> </array> <key>Style</key> <dict> @@ -376,8 +376,8 @@ <integer>45</integer> <key>Points</key> <array> - <string>{135.5000000003933, 71.991527195308294}</string> - <string>{183.50000003815074, 71.983552790886364}</string> + <string>{135.50000000000003, 71.997751477910086}</string> + <string>{183.50000000000406, 71.995635221825466}</string> </array> <key>Style</key> <dict> @@ -725,7 +725,7 @@ <key>MasterSheets</key> <array/> <key>ModificationDate</key> - <string>2014-01-25 01:20:31 +0000</string> + <string>2014-01-29 04:54:54 +0000</string> <key>Modifier</key> <string>MasaKoha</string> <key>NotesVisible</key> @@ -801,7 +801,7 @@ <key>ExpandedCanvases</key> <array/> <key>Frame</key> - <string>{{733, -41}, {694, 878}}</string> + <string>{{768, 313}, {694, 878}}</string> <key>ListView</key> <true/> <key>OutlineWidth</key>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/slide/OUTLINE Wed Jan 29 16:12:51 2014 +0900 @@ -0,0 +1,11 @@ +title:Cerium による並列処理向け I/O の設計と実装 +研究目的・研究背景 +Cerium とは +Cerium での Task の生成法 +並列処理向け I/O でしたこと +mmap +divide read +mmap vs divide read (IO_0 追加前) +Cerium の改良 (IO_0 の追加) +divide read (IO_0 追加前 vs IO_0 追加後) +mmap vs divide read (IO_0 追加後)
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/slide/example.html Wed Jan 29 16:12:51 2014 +0900 @@ -0,0 +1,337 @@ +<!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 による並列処理向け I/O の設計と実装</h1> + <h3 class="title">Masataka Kohagura 12th, February</h3> + <div align="right">担当教官 : 河野 真治</div> + </article> + <article > + <h3>研究背景と目的</h3> + <p> + 近年のCPUのほとんどはマルチコアであり、それらの性能を引き出すためには並列プログラミングが必須となっている。そこで当研究室では Cerium Library の開発を行い、提供することによって並列プログラミングを容易にしている。 + </p> + <p> + 先行研究では Task の並列化によって、プログラム全体の処理速度は向上しているが、ファイル読み込み等の I/O に対して並列に Task が動作するようにはされていない。 + </p> + <p> + 本研究では I/O と Task の並列化の設計、実装によってプログラム全体の処理速度、処理効率を上げていく。 + </p> + </article> + + <article> + <h3>Cerium とは</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> + <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>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/slide/images/createTask.graffle Wed Jan 29 16:12:51 2014 +0900 @@ -0,0 +1,1568 @@ +<?xml version="1.0" encoding="UTF-8"?> +<!DOCTYPE plist PUBLIC "-//Apple//DTD PLIST 1.0//EN" "http://www.apple.com/DTDs/PropertyList-1.0.dtd"> +<plist version="1.0"> +<dict> + <key>ActiveLayerIndex</key> + <integer>0</integer> + <key>ApplicationVersion</key> + <array> + <string>com.omnigroup.OmniGraffle</string> + <string>139.16.0.171715</string> + </array> + <key>AutoAdjust</key> + <true/> + <key>BackgroundGraphic</key> + <dict> + <key>Bounds</key> + <string>{{0, 0}, {1118, 783}}</string> + <key>Class</key> + <string>SolidGraphic</string> + <key>ID</key> + <integer>2</integer> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>stroke</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + </dict> + <key>BaseZoom</key> + <integer>0</integer> + <key>CanvasOrigin</key> + <string>{0, 0}</string> + <key>ColumnAlign</key> + <integer>1</integer> + <key>ColumnSpacing</key> + <real>36</real> + <key>CreationDate</key> + <string>2013-02-18 07:20:36 +0000</string> + <key>Creator</key> + <string>yuhi</string> + <key>DisplayScale</key> + <string>1 0/72 in = 1.0000 in</string> + <key>GraphDocumentVersion</key> + <integer>8</integer> + <key>GraphicsList</key> + <array> + <dict> + <key>Bounds</key> + <string>{{324.209228515625, 389.66677856445312}, {66.0860595703125, 42.153620024593849}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>73</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 CPU}</string> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>71</integer> + </dict> + <key>ID</key> + <integer>72</integer> + <key>Points</key> + <array> + <string>{595.02079148168673, 357.51299844585151}</string> + <string>{595.04302978515625, 386.00642283911986}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>66</integer> + <key>Info</key> + <integer>1</integer> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{562, 386.00642283911986}, {66.0860595703125, 42.153620024593849}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>71</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>stroke</key> + <dict> + <key>Color</key> + <dict> + <key>b</key> + <string>0</string> + <key>g</key> + <string>0</string> + <key>r</key> + <string>1</string> + </dict> + <key>Width</key> + <real>3</real> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 Gpu}</string> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>62</integer> + </dict> + <key>ID</key> + <integer>63</integer> + <key>Points</key> + <array> + <string>{174.28841400146496, 264.65901184082031}</string> + <string>{174.28841484917552, 316.58523266713962}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>56</integer> + <key>Info</key> + <integer>1</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>56</integer> + <key>Info</key> + <integer>2</integer> + </dict> + <key>ID</key> + <integer>57</integer> + <key>Points</key> + <array> + <string>{358.33457709779111, 91.464589235127278}</string> + <string>{174.28841400146496, 207.20574951171875}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>3</integer> + <key>Info</key> + <integer>1</integer> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{95.066192626953352, 316.58523266713962}, {158.44444444444431, 48.883853912353487}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>62</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 FifoScheduler}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{99.288414001464957, 207.20574951171875}, {150, 57.453262329101562}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>56</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 FifoManager}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{327.37777716454013, 314.68888982497248}, {144.65531914893614, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>69</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 Scheduler}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{313.2267133347533, 303.87869152468977}, {144.65531914893614, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>68</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 Scheduler}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{299.07564950496595, 293.06849322440587}, {144.65531914893614, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>70</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 Scheduler}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{88.331925542304646, 34.61473087818672}, {166.66808510638282, 34.832861189801683}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>48</integer> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>fill</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>stroke</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Pad</key> + <integer>0</integer> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 2:spawn()}</string> + <key>VerticalPad</key> + <integer>0</integer> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{77, 15.999999999999776}, {231.1340425531917, 34.832861189801683}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>47</integer> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>fill</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>stroke</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Pad</key> + <integer>0</integer> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 1:createtask()}</string> + <key>VerticalPad</key> + <integer>0</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>46</integer> + <key>Points</key> + <array> + <string>{185.98190489424053, 83.05665722379581}</string> + <string>{249.02658574530432, 54.430594900849705}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>FilledArrow</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>31</integer> + <key>Points</key> + <array> + <string>{187.84360702189986, 96.464589235127377}</string> + <string>{256.88828787296353, 62.838526912181194}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>FilledArrow</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>66</integer> + </dict> + <key>ID</key> + <integer>67</integer> + <key>Points</key> + <array> + <string>{595.24798993049842, 263.65901111678431}</string> + <string>{595.02079148168673, 308.62914453349799}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>60</integer> + <key>Info</key> + <integer>1</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>64</integer> + </dict> + <key>ID</key> + <integer>65</integer> + <key>Points</key> + <array> + <string>{357.25224524964699, 264.65901111678431}</string> + <string>{357.25224524964699, 282.25829492412299}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>58</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>60</integer> + <key>Info</key> + <integer>2</integer> + </dict> + <key>ID</key> + <integer>61</integer> + <key>Points</key> + <array> + <string>{358.33457709779111, 91.464589235127278}</string> + <string>{595.24798993049842, 207.20575332641602}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>57</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>58</integer> + <key>Info</key> + <integer>2</integer> + </dict> + <key>ID</key> + <integer>59</integer> + <key>Points</key> + <array> + <string>{358.33457709779111, 91.464589235127278}</string> + <string>{357.25224524964699, 208.20575332641602}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>3</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>5</integer> + <key>Info</key> + <integer>4</integer> + </dict> + <key>ID</key> + <integer>41</integer> + <key>Points</key> + <array> + <string>{438.00000525941221, 63.237960339943157}</string> + <string>{492.24571622483239, 118.49858403408452}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>3</integer> + <key>Info</key> + <integer>3</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>4</integer> + <key>Info</key> + <integer>4</integer> + </dict> + <key>ID</key> + <integer>40</integer> + <key>Points</key> + <array> + <string>{438.00000525941221, 63.237960339943157}</string> + <string>{492.24571622483239, 63.237962220275321}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>3</integer> + <key>Info</key> + <integer>3</integer> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>3</integer> + <key>Info</key> + <integer>4</integer> + </dict> + <key>ID</key> + <integer>39</integer> + <key>Points</key> + <array> + <string>{166.78723404255317, 119.69121813031163}</string> + <string>{278.66914893617002, 63.237960339943157}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>1</integer> + <key>Info</key> + <integer>3</integer> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{511.68674892849526, 308.62914453349799}, {166.66808510638282, 48.883853912353487}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>66</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>stroke</key> + <dict> + <key>Color</key> + <dict> + <key>b</key> + <string>0</string> + <key>g</key> + <string>0</string> + <key>r</key> + <string>1</string> + </dict> + <key>Width</key> + <real>3</real> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 GpuScheduler}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{284.92458567517889, 282.25829492412299}, {144.65531914893614, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>64</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 Scheduler}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{522.92033035603038, 207.20575332641602}, {144.65531914893614, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>60</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + <key>stroke</key> + <dict> + <key>Color</key> + <dict> + <key>b</key> + <string>0</string> + <key>g</key> + <string>0</string> + <key>r</key> + <string>1</string> + </dict> + <key>Width</key> + <real>3</real> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 GpuThreads}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{284.92458567517889, 208.20575332641602}, {144.65531914893614, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>58</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 CpuThreads}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{492.24571622483245, 96.878187433517922}, {166.66808510638288, 43.240793201133194}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>5</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 WaitTaskList}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{492.24571622483245, 41.617565619708728}, {166.66808510638288, 43.240793201133194}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>4</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;\f1\fnil\fcharset128 HiraKakuProN-W3;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 Ac +\f1 t +\f0 iveTaskList}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{278.66914893617002, 35.011331444759023}, {159.33085632324219, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>3</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 TaskManager}</string> + </dict> + </dict> + <dict> + <key>Bounds</key> + <string>{{41, 91.464589235127491}, {125.78723404255315, 56.453257790368269}}</string> + <key>Class</key> + <string>ShapedGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>ID</key> + <integer>1</integer> + <key>Magnets</key> + <array> + <string>{0, 1}</string> + <string>{0, -1}</string> + <string>{1, 0}</string> + <string>{-1, 0}</string> + <string>{1, 1}</string> + <string>{1, -1}</string> + <string>{-1, 1}</string> + <string>{-1, -1}</string> + </array> + <key>Shape</key> + <string>Rectangle</string> + <key>Style</key> + <dict> + <key>shadow</key> + <dict> + <key>Draws</key> + <string>NO</string> + </dict> + </dict> + <key>Text</key> + <dict> + <key>Text</key> + <string>{\rtf1\ansi\ansicpg1252\cocoartf1187\cocoasubrtf340 +\cocoascreenfonts1{\fonttbl\f0\fswiss\fcharset0 Helvetica;} +{\colortbl;\red255\green255\blue255;} +\pard\tx560\tx1120\tx1680\tx2240\tx2800\tx3360\tx3920\tx4480\tx5040\tx5600\tx6160\tx6720\pardirnatural\qc + +\f0\fs48 \cf0 User Task}</string> + </dict> + </dict> + <dict> + <key>Class</key> + <string>LineGraphic</string> + <key>FontInfo</key> + <dict> + <key>Font</key> + <string>Helvetica</string> + <key>Size</key> + <real>12</real> + </dict> + <key>Head</key> + <dict> + <key>ID</key> + <integer>73</integer> + </dict> + <key>ID</key> + <integer>74</integer> + <key>Points</key> + <array> + <string>{357.25225830078125, 339.21155270420547}</string> + <string>{357.25225830078125, 389.66677856445312}</string> + </array> + <key>Style</key> + <dict> + <key>stroke</key> + <dict> + <key>HeadArrow</key> + <string>0</string> + <key>Legacy</key> + <true/> + <key>TailArrow</key> + <string>0</string> + </dict> + </dict> + <key>Tail</key> + <dict> + <key>ID</key> + <integer>64</integer> + </dict> + </dict> + </array> + <key>GridInfo</key> + <dict/> + <key>GuidesLocked</key> + <string>NO</string> + <key>GuidesVisible</key> + <string>YES</string> + <key>HPages</key> + <integer>2</integer> + <key>ImageCounter</key> + <integer>1</integer> + <key>KeepToScale</key> + <false/> + <key>Layers</key> + <array> + <dict> + <key>Lock</key> + <string>NO</string> + <key>Name</key> + <string>Layer 1</string> + <key>Print</key> + <string>YES</string> + <key>View</key> + <string>YES</string> + </dict> + </array> + <key>LayoutInfo</key> + <dict> + <key>Animate</key> + <string>NO</string> + <key>circoMinDist</key> + <real>18</real> + <key>circoSeparation</key> + <real>0.0</real> + <key>layoutEngine</key> + <string>dot</string> + <key>neatoSeparation</key> + <real>0.0</real> + <key>twopiSeparation</key> + <real>0.0</real> + </dict> + <key>LinksVisible</key> + <string>NO</string> + <key>MagnetsVisible</key> + <string>NO</string> + <key>MasterSheets</key> + <array/> + <key>ModificationDate</key> + <string>2013-02-20 10:08:53 +0000</string> + <key>Modifier</key> + <string>yuhi</string> + <key>NotesVisible</key> + <string>NO</string> + <key>Orientation</key> + <integer>2</integer> + <key>OriginVisible</key> + <string>NO</string> + <key>PageBreaks</key> + <string>YES</string> + <key>PrintInfo</key> + <dict> + <key>NSBottomMargin</key> + <array> + <string>float</string> + <string>41</string> + </array> + <key>NSHorizonalPagination</key> + <array> + <string>coded</string> + <string>BAtzdHJlYW10eXBlZIHoA4QBQISEhAhOU051bWJlcgCEhAdOU1ZhbHVlAISECE5TT2JqZWN0AIWEASqEhAFxlwCG</string> + </array> + <key>NSLeftMargin</key> + <array> + <string>float</string> + <string>18</string> + </array> + <key>NSPaperSize</key> + <array> + <string>size</string> + <string>{595, 842}</string> + </array> + <key>NSPrintReverseOrientation</key> + <array> + <string>int</string> + <string>0</string> + </array> + <key>NSRightMargin</key> + <array> + <string>float</string> + <string>18</string> + </array> + <key>NSTopMargin</key> + <array> + <string>float</string> + <string>18</string> + </array> + </dict> + <key>PrintOnePage</key> + <false/> + <key>ReadOnly</key> + <string>NO</string> + <key>RowAlign</key> + <integer>1</integer> + <key>RowSpacing</key> + <real>36</real> + <key>SheetTitle</key> + <string>Canvas 1</string> + <key>SmartAlignmentGuidesActive</key> + <string>YES</string> + <key>SmartDistanceGuidesActive</key> + <string>YES</string> + <key>UniqueID</key> + <integer>1</integer> + <key>UseEntirePage</key> + <false/> + <key>VPages</key> + <integer>1</integer> + <key>WindowInfo</key> + <dict> + <key>CurrentSheet</key> + <integer>0</integer> + <key>ExpandedCanvases</key> + <array> + <dict> + <key>name</key> + <string>Canvas 1</string> + </dict> + </array> + <key>Frame</key> + <string>{{202, 51}, {991, 922}}</string> + <key>ListView</key> + <true/> + <key>OutlineWidth</key> + <integer>142</integer> + <key>RightSidebar</key> + <false/> + <key>ShowRuler</key> + <true/> + <key>Sidebar</key> + <true/> + <key>SidebarWidth</key> + <integer>120</integer> + <key>VisibleRegion</key> + <string>{{29, 0}, {856, 783}}</string> + <key>Zoom</key> + <real>1</real> + <key>ZoomValues</key> + <array> + <array> + <string>Canvas 1</string> + <real>1</real> + <real>1</real> + </array> + </array> + </dict> +</dict> +</plist>
--- a/slide/index.html Tue Jan 28 18:17:39 2014 +0900 +++ b/slide/index.html Wed Jan 29 16:12:51 2014 +0900 @@ -3,7 +3,7 @@ <html> <head> <title>Presentation</title> - + <meta charset='utf-8'> <script src='./slides.js'></script> @@ -41,11 +41,11 @@ </article> <article> - <h3>Cerium とは</h3> + <h3>Cerium の流れ</h3> <table border="0" cellpadding="0" cellspacing="0"> <tbody> <tr> - <td><img src='images/createtask.png' style="height:350px"></td> + <td><img src='images/cerium.png' style="height:350px"></td> <td> <ol> <li>Taskを生成</li> @@ -57,281 +57,69 @@ </tr> </tbody> </table> - <p> - CpuThreads、Schedulerに対応させる形でGpuThreadsとGpuSchedulerを作成した - </p> - </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> + <h3>Cerium Task の生成(1)</h3> + <br> + <h3 class="yellow">main.cc の記述</h3> + <pre> +// Task の宣言 +HTaskPtr multiply = manager->create_task(MULTIPLY_TASK); +// Task を実行する デバイスの設定 +multiply->set_cpu(SPE_ANY); +// Task に入力データのアドレスを追加 +multiply->set_inData(0, i_data1, sizeof(float)*length); +multiply->set_inData(1, i_data2, sizeof(float)*length); +// Task に出力データのアドレスを追加 +multiply->set_outData(0, o_data1, sizeof(float)*length); +// Task へ値を1つだけ渡す +multiply->set_param(0,(long)length); +// Task を TaskList に set する +multiply->spawn(); </pre> </article> <article> - <h3>Cerium OpenCL API比較</h3> - <img src='images/api.png' style="height:500px"> + <h3>Cerium Task の生成(2)</h3> + <br> + <h3 class="yellow">Task の記述</h3> + <pre> +static int +multiply(SchedTask *s,void *rbuf, void *wbuf) +{ + // 登録した inData を取得 + float indata1=(float*)s->get_input(rbuf,0); + float indata2=(float*)s->get_input(rbuf,1); + // 登録した outData を取得 + float outdata=(float*)s->get_output(wbuf,0); + // 登録した param を取得 + long length=(long)s->get_param(0); + for (int i=0;i < length;i++) { + outdata[i]=indata1[i]*indata2[i]; + } + return 0; +} </pre> </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> + <h3>並列処理向け I/O の 設計と実装</h3> + <br> + <ul> + <li>mmap の仕様</li> + <li>divide read の実装</li> + <li>Cerium の改良(IO_0 の追加)</li> + </ul> </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> + <h3>mmap での I/O の実装</h3> + <br> </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> + <h3>mmap の仕様</h3> + <br> </article> - <!--- <img src='images/flow_chart.jpg' width="300" height="500"> --> + </body> </html>