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