Mercurial > hg > Papers > 2018 > parusu-master
annotate slide/slide.md @ 94:e71d4ea8a616
Fix slide
| author | Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Tue, 13 Feb 2018 02:23:28 +0900 |
| parents | 0050ad0b4dd7 |
| children | 14bd1b466bcc |
| rev | line source |
|---|---|
| 77 | 1 title: Gears OS の並列処理 |
| 2 author: 伊波 立樹 | |
| 3 profile: 琉球大学理工学研究科 河野研 | |
| 4 lang: Japanese | |
| 5 code-engine: coderay | |
| 82 | 6 |
| 89 | 7 ## 並列処理の重要性 |
| 83 | 8 - 並列処理は現在主流のマルチコアCPU の性能を発揮するには重要なものになっている |
| 9 - しかし、並列処理のチューニングや信頼性を保証するのは難しい | |
| 89 | 10 - 共通資源の競合などの非決定的な実行が発生するため、従来のテストやデバッグではテストしきれない部分が残ってしまう |
| 83 | 11 - GPU などのアーキテクチャに合わせた並列プログラミングの記述 |
| 12 | |
| 13 ## Gears OS | |
| 87 | 14 - 本研究室では処理の単位を Code Gear、 データの単位を Data Gear を用いて 信頼性が高い並列処理を行う Gears OS を開発している |
| 15 - 並列処理の Task を Code Gear と実行するときに必要な Input Data Gear と出力するための Output Data Gear の組で表現される | |
| 16 - 計算をノーマルレベルとメタレベルに階層化、 信頼性と拡張性をメタレベルで保証する | |
| 83 | 17 - 並列処理の信頼性を通常の計算(ノーマルレベル) に保証 |
| 89 | 18 - CPU、GPU などの実行環境の切り替え、データ拡張等を提供 |
| 83 | 19 |
| 20 ## Gears OS | |
| 84 | 21 - 本研究ではGears OS の並列処理機構、並列処理構文(par goto)の実装、Gears OS を実装するにつれて必要なったモジュール化の導入を行う |
| 87 | 22 - また、並列処理を行う例題を用いて評価、 OpenMP、 Go 言語との比較を行う |
| 84 | 23 |
| 89 | 24 ## Code Gear/Data Gear |
| 87 | 25 - Gears OS は Code Gear、 Data Gear という単位で構成される |
| 77 | 26 - Code Gear はプログラムの処理そのものを表す |
| 27 - Data Gear はデータそのものを表す | |
| 82 | 28 - Code Gear は必要な Input Data Gear が揃ったら実行し、Output Data Gear を生成する |
| 29 - Code Gear と Input / Output Data Gear の対応から依存関係を解決し、Input Data Gear が揃った Code Gear の並列実行を行う | |
| 77 | 30 |
| 31 <div style="text-align: center;"> | |
| 94 | 32 <img src="./images/codegear-datagear-dependency.svg" alt="message" width="600"> |
| 77 | 33 </div> |
| 34 | |
| 87 | 35 ## メタ計算 |
| 36 - メタ計算 は通常の計算を実行するための計算 | |
| 37 - 信頼性の確保やメモリ管理、スレッド管理、CPU、GPU の資源管理等 | |
| 38 - Gears OS のメタ計算は通常の計算とは別の階層のメタレベルで行われる | |
| 39 - メタレベルは Code/Data Gear に対応して Meta Code/Data Gear で行われる | |
| 40 | |
| 41 ## Meta Gear | |
| 42 - メタ計算 は Code Gear の接続間に行われる | |
| 43 - この Gear を Meta Code/Data Gearと呼ぶ | |
| 44 - Meta Code Gear は メタ計算 のプログラム部分 | |
| 45 - Meta Data Gear は Meta Code Gear で管理されるデータ部分 | |
| 46 - Gears OS は通常の Code/Data Gear から Meta Code/Data Gear 部分は見えないように実装を行う | |
| 47 | |
| 48 <div style="text-align: center;"> | |
| 49 <img src="./images/meta_cg_dg.svg" alt="message" width="850"> | |
| 50 </div> | |
| 51 | |
| 77 | 52 ## Continuation based C |
| 53 - Gears OS の実装は本研究室で開発している Continuation based C(CbC) を用いる | |
| 87 | 54 - CbC は Code Gear を用いて記述する事を基本とする |
| 77 | 55 |
| 56 ## Continuation based C | |
| 88 | 57 - CbC の Code Gear の定義は **__code CG名** で行う |
| 58 - Code Gear 間は **goto CG名** で移動する。この移動を継続と呼ぶ | |
| 77 | 59 - Code Gear の継続は C の関数呼び出しとは異なり、戻り値を持たないためスタックの変更を行わない |
| 60 - このような環境を持たない継続を軽量継続と呼ぶ | |
| 88 | 61 - 下記のコードでは Code Gear を2つ定義し、 cg0 から cg1 への継続を示している |
| 77 | 62 |
| 63 ``` c | |
| 64 __code cg0(int a, int b) { | |
| 65 goto cg1(a+b); | |
| 66 } | |
| 67 | |
| 68 __code cg1(int c) { | |
| 69 goto cg2(c); | |
| 70 } | |
| 71 ``` | |
| 72 | |
| 73 ## Data Gear の表現 | |
| 74 - Data Gear は構造体を用いて定義する | |
| 75 - メタ計算では任意の Data Gear を一律に扱うため、全ての Data Gear は共用体の中で定義される | |
| 76 - Data Gear のメモリに確保する際のサイズ情報はこの型から決定する | |
| 77 | |
| 78 ``` c | |
| 79 /* data Gear define */ | |
| 80 union Data { | |
| 81 struct Timer { | |
| 82 union Data* timer; | |
| 83 enum Code start; | |
| 84 enum Code end; | |
| 85 enum Code next; | |
| 86 } Timer; | |
| 87 struct TimerImpl { | |
| 88 double time; | |
| 89 } TimerImpl; | |
| 90 .... | |
| 91 }; | |
| 92 ``` | |
| 93 | |
| 94 | 94 ## Context |
| 95 - Context は従来のOS のスレッドやプロセスに対応し、以下の要素をもっている Meta Data Gear | |
| 96 - Data Gear を確保するためのメモリ空間 | |
| 97 - Code Gear の名前と関数ポインタとの対応表 | |
| 98 - Code Gear は番号(enum)で指定する | |
| 99 - Code Gear が参照する Data Gear へのポインタ | |
| 100 - Code Gear と同じく Data Gear も番号で指定する | |
| 101 - 並列実行用の Task 情報 | |
| 102 - Data Gear の型情報 | |
| 103 - Gears OS ではメタ計算で Context を経由して Code/Data Gear にアクセスする | |
| 104 | |
| 88 | 105 ## stub Code Gear |
| 77 | 106 - Data Gear にアクセスするにはContext を経由する |
| 107 - だが、通常の Code Gear では Meta Data Gear である Context の参照は避ける必要がある | |
| 89 | 108 - Gears OS ではメタレベルで通常の Code Gear で必要な Data Gear を Context から取り出す処理を行う stub Code Gear を用意している |
| 77 | 109 |
| 94 | 110 <div style="text-align: center;"> |
| 111 <img src="./images/contextContinuation.svg" alt="message" width="600"> | |
| 112 </div> | |
| 87 | 113 |
| 77 | 114 |
| 115 ## Context での stub Code Gear の記述の問題点 | |
| 94 | 116 - Gears OS を実装するに連れて、stub Code Gear の記述が煩雑になる場所がでてきた |
| 77 | 117 - そのため Gears OS のモジュール化する仕組みとして **Interface** を導入した |
| 118 | |
| 119 ## Interface | |
| 120 - Interface はある Data Gear と それに対する操作(API) を行う Code Gear の集合を表現する Meta Data Gear | |
| 88 | 121 - stub Code Gear はInteface を実装した Code Gear で決まった形になるため、自動生成が可能である |
| 87 | 122 - Interface を導入することで、 Stack や Queue などのデータ構造を仕様と実装に分けて記述することが出来る |
| 77 | 123 - Interface は Java のインターフェース、 Haskell の型クラスに対応する |
| 82 | 124 |
| 125 ## Interface の定義 | |
| 85 | 126 - Interface の定義には以下の内容を定義する |
| 94 | 127 - 操作(API)の引数群の型 |
| 128 - 操作(API)自体のCode Gear の型 | |
| 85 | 129 |
| 130 ``` c | |
| 131 typedef struct Queue<Impl>{ | |
| 132 // Data Gear parameter | |
| 133 union Data* queue; | |
| 134 union Data* data; | |
| 135 __code next(...); | |
| 136 __code whenEmpty(...); | |
| 137 | |
| 138 // Code Gear | |
| 139 __code clear(Impl* queue, __code next(...)); | |
| 140 __code put(Impl* queue, union Data* data, __code next(...)); | |
| 141 __code take(Impl* queue, __code next(union Data*, ...)); | |
| 142 __code isEmpty(Impl* queue, __code next(...), __code whenEmpty(...)); | |
| 143 } Queue; | |
| 144 ``` | |
| 82 | 145 |
| 146 ## Interface の実装 | |
| 85 | 147 - Interface には複数の実装を行うことが出来る |
| 94 | 148 - 実装した Code Gear の番号を Interface の定義に代入することで実装を行う |
| 85 | 149 - 代入する Code Gear を入れ替えることで別の実装を表現する |
| 86 | 150 - 実装した Data Gear の生成は関数呼び出しで行われ、外から見るとInterface の型で扱われる |
| 85 | 151 |
| 152 ``` | |
| 153 Queue* createSingleLinkedQueue(struct Context* context) { | |
| 154 struct Queue* queue = new Queue(); // Allocate Queue interface | |
| 155 struct SingleLinkedQueue* singleLinkedQueue = new SingleLinkedQueue(); // Allocate Queue implement | |
| 156 queue->queue = (union Data*)singleLinkedQueue; | |
| 157 singleLinkedQueue->top = new Element(); | |
| 158 singleLinkedQueue->last = singleLinkedQueue->top; | |
| 159 queue->clear = C_clearSingleLinkedQueue; | |
| 160 queue->put = C_putSingleLinkedQueue; | |
| 161 queue->take = C_takeSingleLinkedQueue; | |
| 162 queue->isEmpty = C_isEmptySingleLinkedQueue; | |
| 163 return queue; | |
| 164 } | |
| 165 ``` | |
| 166 | |
| 167 ## Interface の実装例 | |
| 168 - SingleLinkedQueue の put 実装 | |
| 88 | 169 - 引数は Queue Interface の put 定義にあわせる |
| 86 | 170 - 第1引数は 実装対象の Data Gear の型になる |
| 171 - 第3引数の(...) は Output Data Gear を記述する | |
| 172 - ... は可変長引数のような扱いで、 継続先の Code Gear が複数の値をInput Data Gear とする可能性がある | |
| 85 | 173 |
| 174 ``` c | |
| 175 __code putSingleLinkedQueue(struct SingleLinkedQueue* queue, union Data* data, __code next(...)) { | |
| 176 Element* element = new Element(); | |
| 177 element->data = data; | |
| 178 element->next = NULL; | |
| 179 queue->last->next = element; | |
| 180 queue->last = element; | |
| 181 goto next(...); | |
| 182 } | |
| 183 ``` | |
| 82 | 184 |
| 185 ## Interface を利用した Code Gear の呼び出し | |
| 86 | 186 - Interface を利用した Code Gear への継続は `goto interface->method` で行われる |
| 89 | 187 - ここでの **interface** は Interfaceの型で包んだData Gear、 **method** は実装した Code Gear に対応する |
| 86 | 188 |
| 189 ``` | |
| 94 | 190 __code code1() { |
| 86 | 191 Queue* queue = createSingleLinkedQueue(context); |
| 192 Node* node = new Node(); | |
| 193 node->color = Red; | |
| 194 goto queue->put(node, queueTest2); | |
| 195 } | |
| 196 ``` | |
| 197 | |
| 198 ## Interface を利用した Code Gear の呼び出し(スクリプト変換後) | |
| 88 | 199 - Interface を利用した Code Gear の継続はスクリプトによって変換される |
| 200 - 変換後は Context を参照するため、メタレベルの記述になる | |
| 86 | 201 - Gearef マクロは Context から Interface の引数格納用の Data Gear を取り出す |
| 202 - この Data Gear は Context を初期化した際に特別に生成され、型は Interface と同じになる | |
| 88 | 203 - 呼び出すCode Gear の引数情報に合わせて引数に格納し、 実装された Code Gear へ継続する |
| 86 | 204 |
| 205 ``` | |
| 206 __code code1(struct Context *context) { | |
| 207 Queue* queue = createSingleLinkedQueue(context); | |
| 208 Node* node = &ALLOCATE(context, Node)->Node; | |
| 209 node->color = Red; | |
| 210 Gearef(context, Queue)->queue = (union Data*) queue; | |
| 211 Gearef(context, Queue)->data = (union Data*) node; | |
| 212 Gearef(context, Queue)->next = C_queueTest2; | |
| 213 goto meta(context, queue->put); | |
| 214 } | |
| 215 ``` | |
| 216 | |
| 217 ## Interface での stub Code Gear | |
| 88 | 218 - メタ計算で格納された引数は stub Code Gear で Code Gear に渡される |
| 219 - Interface を実装した Code Gear は Interface の定義から stub Code Gear の自動生成が可能である | |
| 86 | 220 |
| 221 ``` c | |
| 222 __code putSingleLinkedQueue(struct Context *context,struct SingleLinkedQueue* queue, union Data* data, enum Code next) { | |
| 223 Element* element = &ALLOCATE(context, Element)->Element; | |
| 224 element->data = data; | |
| 225 element->next = NULL; | |
| 226 queue->last->next = element; | |
| 227 queue->last = element; | |
| 228 goto meta(context, next); | |
| 229 } | |
| 230 | |
| 231 // generated by script | |
| 232 __code putSingleLinkedQueue_stub(struct Context* context) { | |
| 233 SingleLinkedQueue* queue = (SingleLinkedQueue*)GearImpl(context, Queue, queue); | |
| 234 Data* data = Gearef(context, Queue)->data; | |
| 235 enum Code next = Gearef(context, Queue)->next; | |
| 236 goto putSingleLinkedQueue(context, queue, data, next); | |
| 237 } | |
| 238 ``` | |
| 239 | |
| 82 | 240 ## 並列処理の構成 |
| 94 | 241 - 今回は並列処理機構である |
| 92 | 242 - Task |
| 86 | 243 - TaskManager |
| 244 - Worker の生成、依存関係を解決したTask を Worker に送信する | |
| 77 | 245 - Worker |
| 92 | 246 - SynchronizedQueue から Task を一つずつ取得し、実行する |
| 247 - Worker は CPU、 GPU の対応した数分生成する | |
| 248 - Worker 毎に POSIX Therad などのスレッドを生成し、 Code Gear を実行する | |
| 88 | 249 - SynchronizedQueue |
| 250 - マルチスレッド 環境でもデータの一貫性を保証する Queue | |
| 94 | 251 - をInterface で実装した |
| 77 | 252 |
| 92 | 253 ## Task |
| 254 - Gears OS では Context が並列実行の Task に相当する | |
| 88 | 255 - Context は Task用の情報として以下の情報をもっている |
| 86 | 256 - 実行する Code Gear |
| 257 - Input/Output Data Gear の格納場所 | |
| 258 - 待っている Input Data Gear の数 | |
| 77 | 259 |
| 82 | 260 ## TaskManger |
| 92 | 261 - Worker を作成、終了処理を行う |
| 82 | 262 - 依存関係を解決した Task を各 Worker の Queue に送信する |
| 77 | 263 |
|
90
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
264 <div> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
265 <div style="float: left;"> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
266 <img src="./images/sendTask.svg" alt="message" width="600"> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
267 </div> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
268 <div style="float: left; font-size=100%;"> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
269 <ol> |
| 92 | 270 <li value="0">Task を Input Data Gear として</li> |
| 271 TaskManager の spawn を呼び出す | |
| 94 | 272 <li value="1">Task内の idgCountをチェックする</li> |
| 273 idgCount は待っているInput Data Gear のカウンタである | |
| 274 そのため、カウンタが0の場合 Input Data Gear が揃っていることになる | |
| 92 | 275 <li value="2">揃っている場合、 Worker の Queue に</li> |
| 276 Task を送信する | |
|
90
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
277 </ol> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
278 </div> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
279 <div style="clear: both;"></div> |
| 82 | 280 </div> |
| 77 | 281 |
| 282 ## Worker | |
| 92 | 283 - 初期化時に Worker 用の Context を生成する |
| 94 | 284 - TaskManager から送信された Task を一つずつ取得して実行する |
| 86 | 285 |
| 286 <div> | |
| 287 <div style="float: left;"> | |
| 288 <img src="./images/workerRun.svg" alt="message" width="600"> | |
| 289 </div> | |
| 290 <div style="float: left; font-size=100%;"> | |
| 291 <ol> | |
| 292 <li>Worker は Queue から Task を取得する</li> | |
|
90
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
293 <li>Worker の Context から</li> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
294 Task の Context へ入れ替える |
| 92 | 295 <li>Task に設定されている Code Gear を実行</li> |
| 86 | 296 <li>Task の Output Data Gear の書き出し</li> |
|
90
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
297 <li>Task Context から</li> |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
298 Worker の Context へ入れ替える |
| 86 | 299 <li>Worker は再び Queue から Task を取得する</li> |
| 300 </ol> | |
| 301 </div> | |
| 88 | 302 <div style="clear: both;"></div> |
| 86 | 303 </div> |
| 304 | |
| 305 ## Synchronized Queue | |
| 306 - Worker で使用される Queue | |
| 92 | 307 - Task を送信するTaskManagerと Task を取得するWorker毎で操作される |
| 94 | 308 - そのためマルチスレッドでのデータの同期処理を行える SynchronizedQueue を実装する |
| 309 - Gears OS では CAS(Check and Set、 Compare and Swap) を使用した実装を行った | |
| 310 - CAS は値を更新する際に更新前の値と実際に保存されているメモリ番地の値を比較し、変化がなければ値を更新する | |
| 311 - メモリ番地の値が変わっているなら、もう一度 CAS を行う | |
| 86 | 312 |
| 94 | 313 <div style="text-align: center;"> |
| 314 <img src="./images/synchronizedQueue.svg" alt="message" width="600"> | |
| 315 </div> | |
| 86 | 316 |
| 317 ## 依存関係の解決 | |
| 88 | 318 - 依存関係の解決は Data Gear がメタレベルで持っている Queue を使用する |
| 92 | 319 - この Queue には Data Gear に依存関係がある Context が格納されている |
| 86 | 320 |
| 88 | 321 <div> |
| 322 <div style="float: left;"> | |
| 323 <img src="./images/dependency.svg" alt="message" width="550"> | |
| 324 </div> | |
| 325 <div style="float: left; font-size=100%;"> | |
| 326 <ol> | |
| 92 | 327 <li>Task に設定されている Code Gear を実行する</li> |
| 88 | 328 <li>Output Data Gear の書き出し処理を行う</li> |
| 329 この際にメタレベルの Queue を参照する | |
| 330 <li>依存関係にある Task を取り出し、 待っている</li> | |
| 331 Data Gearのカウンタをデクリメントする | |
| 332 </ol> | |
| 333 </div> | |
| 334 <div style="clear: both;"></div> | |
| 86 | 335 </div> |
| 336 | |
| 88 | 337 - カウンタの値が0になった実行可能な Task は TaskManager を通して Worker に送信される |
| 338 | |
| 86 | 339 ## 並列構文 |
| 340 - 並列実行の Task の生成は新しく Context を生成し、実行する Code Gear、待ち合わせる Input Data Gear の数、Input/Output Data Gear への参照を設定する | |
| 89 | 341 - この記述を直接書くと Meta Data Gear である Context を直接参照しているため、ノーマルレベルでの記述では好ましくない |
| 342 - Task の設定は煩雑な記述であるが、並列実行されることを除けば通常の CbC の goto 文と同等である | |
| 86 | 343 - そこで Context を直接参照しない並列構文、 **par goto** 構文を新たに考案した |
| 344 - par goto 構文には引数として Input/Output Data Gear等を渡す | |
| 345 - スクリプトによって Code Gear の Input/Output の数を解析する | |
| 346 | |
| 347 ``` c | |
| 348 __code code1(Integer *integer1, Integer * integer2, Integer *output) { | |
| 349 par goto add(integer1, integer2, output, __exit); | |
| 350 goto code2(); | |
| 351 } | |
| 352 ``` | |
| 77 | 353 |
| 86 | 354 ## CUDA への対応 |
| 355 - Gears OS は GPU での実行もサポートする | |
| 89 | 356 - CUDA は GPU を Device、 CPU を Host として定義する |
| 86 | 357 - CUDA は処理の最小の単位を thread とし、それをまとめた block を展開し Device 上で実行されるプログラム(Kernel)を実行する |
| 92 | 358 - GPU で性能を出すためには GPU に合わせた並列プログラミングが必要になる |
| 359 - 今回、CUDA に合わせた並列処理機構を Interface を用いて実装した | |
| 86 | 360 |
| 88 | 361 ## CUDAWorker |
| 362 - CUDA で実行する Task を受け取る Worker | |
| 363 - 初期化の際に CUDA ライブラリの初期化等を行う | |
| 86 | 364 |
| 365 ## CUDAExecutor | |
| 88 | 366 - CUDAExecutor は Executor Interface を実装した以下の Code Gear を持つ |
| 86 | 367 - HostからDevice へのデータの送信(read) |
| 368 - kernel の実行(exec) | |
| 369 - Device から Host へのデータの書き出し(write) | |
| 370 | |
| 371 ``` c | |
| 372 typedef struct Executor<Impl>{ | |
| 373 union Data* Executor; | |
| 374 struct Context* task; | |
| 375 __code next(...); | |
| 376 // method | |
| 377 __code read(Impl* executor, struct Context* task, __code next(...)); | |
| 378 __code exec(Impl* executor, struct Context* task, __code next(...)); | |
| 379 __code write(Impl* executor, struct Context* task, __code next(...)); | |
| 380 } | |
| 381 ``` | |
| 382 | |
| 383 ## CUDABuffer | |
| 88 | 384 - Host、Device 間でデータのやり取りをする際、 Gears OS での Data Gear をDevice 用にマッピングする必要がある |
| 86 | 385 - Device にデータ領域を確保するにはサイズの指定が必要 |
| 386 - Data Gear には Meta Data Gear でデータのサイズを持っている | |
| 94 | 387 - Data Gear の要素の中に Data Gear へのポインタがあるとポインタ分でサイズ計算してしまうため、 GPU では参照できなくなってしまう |
| 86 | 388 - CUDA Buffer ではそのマッピングを行う |
| 389 - このマッピングは Task の stub Code Gear で行われる | |
| 390 | |
| 391 <div style="text-align: center;"> | |
| 91 | 392 <img src="./images/cudaDataArchitecture.svg" alt="message" width="400"> |
| 86 | 393 </div> |
| 394 | |
| 395 ## CUDA での呼び出し | |
| 89 | 396 - Gears OS では Task で実行される Code Gear の stub Code Gear で CUDA による実行を切り替える |
| 86 | 397 - stub Code Gear で CUDABuffer でのマッピング、 実行する kernel の読み込みを行う |
| 398 - stub Code Gear は CUDA で実行する際は CUDAExecutor の Code Gear に継続する | |
| 399 | |
| 400 ## Gears OS の評価 | |
| 401 - 並列処理のタスクの例題として Twice と BitonicSort を実装し、 以下の環境で測定を行った | |
| 402 - CPU 環境 | |
| 403 - Model : Dell PowerEdgeR630 | |
| 404 - Memory : 768GB | |
| 405 - CPU : 2 x 18-Core Intel Xeon 2.30GHz | |
| 406 - GPU 環境 | |
| 407 - GPU : GeForce GTX 1070 | |
| 408 - Cores : 1920 | |
| 409 - ClockSpeed : 1683MHz | |
| 410 - Memory Size : 8GB GDDR5 | |
| 411 - Memory Bandwidth : 256GB/s | |
| 412 | |
| 413 ## Twice | |
| 414 - Twice は与えられた整数配列を2倍にする例題である | |
| 415 - 並列実行の依存関係がなく、並列度が高い課題である | |
| 88 | 416 |
| 417 ## Twice の結果 | |
| 86 | 418 - 要素数 2^27 |
| 419 - CPU での実行時は 2^27 を 2^6 個に分割して Task を生成する | |
| 420 - GPU での実行時は1次元の block 数を 2^15、 block 内の thread 数を 2^10 で展開 | |
| 88 | 421 - 1 CPU と 32 CPU では 約27.1倍の速度向上が見られた |
| 422 - GPU 実行は kernel のみの実行時間は32 CPU に比べて約7.2倍の速度向上、通信時間を含めると 16 CPU より遅い | |
| 423 - 通信時間がオーバーヘッドになっている | |
| 86 | 424 |
| 425 <table border="1" align='center' width='50%'> | |
| 77 | 426 <tbody> |
| 427 <tr> | |
| 86 | 428 <td style="text-align: center;">Processor</td> |
| 429 <td style="text-align: center;">Time(ms)</td> | |
| 430 </tr> | |
| 431 <tr> | |
| 432 <td style="text-align: center;">1 CPU</td> | |
| 433 <td style="text-align: right;">1181.215</td> | |
| 434 </tr> | |
| 435 <tr> | |
| 436 <td style="text-align: center;">2 CPUs</td> | |
| 437 <td style="text-align: right;">627.914</td> | |
| 438 </tr> | |
| 439 <tr> | |
| 440 <td style="text-align: center;">4 CPUs</td> | |
| 441 <td style="text-align: right;">324.059</td> | |
| 442 </tr> | |
| 443 <tr> | |
| 444 <td style="text-align: center;">8 CPUs</td> | |
| 445 <td style="text-align: right;">159.932</td> | |
| 446 </tr> | |
| 447 <tr> | |
| 448 <td style="text-align: center;">16 CPUs</td> | |
| 449 <td style="text-align: right;">85.518</td> | |
| 450 </tr> | |
| 451 <tr> | |
| 452 <td style="text-align: center;">32 CPUs</td> | |
| 453 <td style="text-align: right;">43.496</td> | |
| 454 </tr> | |
| 455 <tr> | |
| 456 <td style="text-align: center;">GPU</td> | |
| 87 | 457 <td style="text-align: right;">127.018</td> |
| 86 | 458 </tr> |
| 459 <tr> | |
| 460 <td style="text-align: center;">GPU(kernel only)</td> | |
| 461 <td style="text-align: right;">6.018</td> | |
| 77 | 462 </tr> |
| 463 </tbody> | |
| 464 </table> | |
| 465 | |
| 82 | 466 ## BitonicSort |
| 86 | 467 - 並列処理向けのソートアルゴリズム |
| 468 - 決まった2点間の要素の入れ替えをステージ毎に並列に実行し、 Output Data Gear として書き出し、次のステージの Code Gear の Input Data Gear とする | |
| 469 | |
| 82 | 470 ## BitonicSort の結果 |
| 88 | 471 - 要素数 2^24 |
| 472 - CPU での実行時は 2^24 を 2^6 個に分割して Task を生成する | |
| 473 - GPU での実行時は1次元の block 数を 2^14、 block 内の thread 数を 2^10 で展開 | |
| 474 - 1 CPU と 32 CPU で約22.12倍の速度向上 | |
| 475 - GPU は通信時間を含めると 8 CPU の約1.16倍、 kernel のみの実行では 32 CPU の約11.48倍になった | |
| 476 - 現在の Gears OS の CUDA 実装では Output Data Gear を書き出す際に一度 GPU から CPU へ kernel の結果の書き出しを行っているため、差がでてしまった | |
| 77 | 477 |
| 478 <table border="1" align='center' width='50%'> | |
| 479 <tbody> | |
| 480 <tr> | |
| 86 | 481 <td style="text-align: center;">Processor</td> |
| 482 <td style="text-align: center;">Time(s)</td> | |
| 77 | 483 </tr> |
| 484 <tr> | |
| 485 <td style="text-align: center;">1 CPU</td> | |
| 86 | 486 <td style="text-align: right;">41.416</td> |
| 77 | 487 </tr> |
| 488 <tr> | |
| 489 <td style="text-align: center;">2 CPUs</td> | |
| 86 | 490 <td style="text-align: right;">23.340</td> |
| 77 | 491 </tr> |
| 492 <tr> | |
| 493 <td style="text-align: center;">4 CPUs</td> | |
| 86 | 494 <td style="text-align: right;">11.952</td> |
| 77 | 495 </tr> |
| 496 <tr> | |
| 497 <td style="text-align: center;">8 CPUs</td> | |
| 86 | 498 <td style="text-align: right;">6.320</td> |
| 499 </tr> | |
| 500 <tr> | |
| 501 <td style="text-align: center;">16 CPUs</td> | |
| 502 <td style="text-align: right;">3.336</td> | |
| 77 | 503 </tr> |
| 504 <tr> | |
| 86 | 505 <td style="text-align: center;">32 CPUs</td> |
| 506 <td style="text-align: right;">1.872</td> | |
| 507 </tr> | |
| 508 <tr> | |
| 509 <td style="text-align: center;">GPU</td> | |
| 510 <td style="text-align: right;">5.420</td> | |
| 511 </tr> | |
| 512 <tr> | |
| 513 <td style="text-align: center;">GPU(kernel only)</td> | |
| 514 <td style="text-align: right;">0.163</td> | |
| 77 | 515 </tr> |
| 516 </tbody> | |
| 517 </table> | |
| 518 | |
| 88 | 519 |
| 520 ## OpenMP との比較 | |
| 521 - OpenMP は C、 C++ のプログラムにアノテーションを付けることで並列化を行う | |
| 522 - データの待ち合わせ処理はバリア等のアノテーションで記述する | |
| 523 - Gears OS は並列処理を par goto 構文、 データの待ち合わせを Code Gear と Input/Ouput Data Gear の関係で行う | |
| 524 | |
| 525 ``` c | |
| 526 #pragma omp parallel for | |
| 527 for(int i = 0; i < length; i++) { | |
| 528 a[i] = a[i] * 2; | |
| 529 } | |
| 530 ``` | |
| 77 | 531 |
| 88 | 532 ## Go 言語との比較 |
| 533 - Go 言語は並列実行を **go funciton(argv)** の構文で行う。 この実行を goroutine と呼ぶ | |
| 534 - データの待ち合わせはチャネルというデータ構造で行う | |
| 535 - チャネルでのデータの送受信は **<-** を使用して行うため、簡潔に書くことが出来る | |
| 536 - しかし、 チャネルは複数の goroutine で共有されるため、データの送信元が推測しづらい | |
| 537 - Gears OS では goroutine は par goto 文とほぼ同等に扱える | |
| 538 - par goto 文では書き出す Data Gear を指定するため、書き出し元が推測しやすい | |
| 539 | |
| 540 ``` go | |
| 541 c := make(chan []int) | |
| 542 for i :=0; i < *split; i++ { | |
| 543 // call goroutine | |
| 544 go twice(list, prefix, i, c); | |
| 545 } | |
| 546 | |
| 547 func twice(list []int, prefix int, index int, c chan []int) { | |
| 548 for i := 0; i < prefix; i++ { | |
| 549 list[prefix*index+i] = list[prefix*index+i] * 2; | |
| 550 } | |
| 551 c <- list | |
| 552 } | |
| 553 ``` | |
| 554 | |
| 92 | 555 ## OpenMP との比較 |
| 556 - OpenMP で Twice を実装し、速度比較を行った | |
| 557 - OpenMP は 1CPU と 32CPU で約10.8倍の速度向上が見られた | |
| 558 - 一方 Gears OS では約27.1倍と台数効果は高くなっている | |
| 559 - しかし、 Gears OS は 1CPU の実行速度が OpenMP に比べて大幅に遅くなっている | |
| 560 | |
| 561 <div style="text-align: center;"> | |
| 562 <img src="./images/vsopenmp.svg" alt="message" width="500"> | |
| 563 </div> | |
| 564 | |
| 88 | 565 ## Go 言語との比較 |
| 566 - Go 言語でも OpenMP と同様に Twice を実装し、速度比較を行った | |
| 567 - Go 言語は 1CPU と 32CPU で約4.33倍の速度向上が見られた | |
| 568 - OpenMP と同様に台数効果自体は Gears OS が高いが、 1CPU での実行時間は Go 言語が大幅に速い | |
| 569 | |
| 570 <div style="text-align: center;"> | |
| 571 <img src="./images/vsgo.svg" alt="message" width="500"> | |
| 572 </div> | |
| 77 | 573 |
| 574 ## まとめ | |
| 86 | 575 - Gears OS の並列処理機構の実装を行った |
| 576 - Interface を導入することで、見通しの良し Gears OS のプログラミングが可能となった | |
| 577 - par goto 構文を導入することで、ノーマルレベルで並列処理の記述が可能になった | |
| 578 - 2つの例題である程度の台数効果が出ることを確認した | |
| 77 | 579 |
| 580 ## 今後の課題 | |
| 86 | 581 - Gears OS の並列処理の信頼性の保証、チューニングを行う |
| 582 - Gears OS では検証とモデル検査をメタレベルで実現することで信頼性を保証する | |
| 88 | 583 - 証明は CbC のプログラムを証明支援系の Agda に対応して行う。 並列処理の信頼性を保証するには SynchronizedQueue の証明を行う必要がある |
| 86 | 584 - モデル検査は CbC で記述された モデル検査器である akasha を使用して行う。 モデル検査の方針としては Code Gear の並列実行を擬似並列で実行し、全ての組合せを列挙する方法で行う |
|
90
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
585 - 現在の CUDA 実装では CPU、GPU 間のデータの通信コストがかかってしまうことが例題からわかった |
|
a9885c038bb6
Add TaskManager describe
Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
parents:
89
diff
changeset
|
586 - Meta Data Gear に Data Gear が CPU、 GPU のどこで所持されているのかを持たせ、 GPU の Data Gear が CPU で必要になったときに始めてデータの通信を行う |
| 86 | 587 - OpenMP、 Goとの比較から、 Gears OS が 1CPU での動作が遅いということがわかった。 |
| 588 - par goto 文を使用する度に Context を生成するため、 ある程度の時間がかかってしまう | |
| 89 | 589 - モデル検査で par goto の Code Gear のフローを解析し、処理がかる場合は Context を生成せずに関数呼出しを行う等の最適化が必要 |