Mercurial > hg > Papers > 2018 > parusu-master

changeset 88:3c127f675c45

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Update slide
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Mon, 12 Feb 2018 12:07:12 +0900
parents 015f9933b245
children bc7d11285a4a
files slide/slide.html slide/slide.md
diffstat 2 files changed, 329 insertions(+), 171 deletions(-) [+]
line wrap: on
line diff
--- a/slide/slide.html	Mon Feb 12 03:16:15 2018 +0900
+++ b/slide/slide.html	Mon Feb 12 12:07:12 2018 +0900
@@ -87,7 +87,7 @@
 <!-- === begin markdown block ===
 
       generated by markdown/1.2.0 on Ruby 2.3.0 (2015-12-25) [x86_64-darwin16]
-                on 2018-02-11 12:53:21 +0900 with Markdown engine kramdown (1.13.2)
+                on 2018-02-12 04:41:22 +0900 with Markdown engine kramdown (1.13.2)
                   using options {}
   -->
 
@@ -110,14 +110,9 @@
 <!-- _S9SLIDE_ -->
 <h2 id="gears-os">Gears OS</h2>
 <ul>
-  <li>本研究室では 処理の単位を Code Gear、 データの単位を Data Gear を用いて 信頼性が高い並列処理を行う Gears OS を開発している</li>
-  <li>Gears OS では Task を Code Gear と実行するときに必要な Input Data Gear と出力するための Output Data Gear の組で表現される
-    <ul>
-      <li>Input Data Gear/Output Data Gear によって依存関係が決定し、それにそって並列実行を行う</li>
-    </ul>
-  </li>
-  <li>Gears OS では計算をノーマルレベルとメタレベルに階層化</li>
-  <li>信頼性と拡張性をメタレベルで保証する
+  <li>本研究室では処理の単位を Code Gear、 データの単位を Data Gear を用いて 信頼性が高い並列処理を行う Gears OS を開発している</li>
+  <li>並列処理の Task を Code Gear と実行するときに必要な Input Data Gear と出力するための Output Data Gear の組で表現される</li>
+  <li>計算をノーマルレベルとメタレベルに階層化、 信頼性と拡張性をメタレベルで保証する
     <ul>
       <li>並列処理の信頼性を通常の計算(ノーマルレベル) に保証</li>
       <li>CPU、GPU などの実行環境の切り替え、データ拡張等のを提供</li>
@@ -132,7 +127,7 @@
 <h2 id="gears-os-1">Gears OS</h2>
 <ul>
   <li>本研究ではGears OS の並列処理機構、並列処理構文(par goto)の実装、Gears OS を実装するにつれて必要なったモジュール化の導入を行う</li>
-  <li>また、並列処理を行う例題を用いて評価、 Open MP、 Go 言語との比較を行う</li>
+  <li>また、並列処理を行う例題を用いて評価、 OpenMP、 Go 言語との比較を行う</li>
 </ul>
 
 
@@ -141,7 +136,7 @@
 <!-- _S9SLIDE_ -->
 <h2 id="code-gear-data-gear">Code Gear、 Data Gear</h2>
 <ul>
-  <li>Gears OS は Code Gear、 Data Gear という Gear で構成される</li>
+  <li>Gears OS は Code Gear、 Data Gear という単位で構成される</li>
   <li>Code Gear はプログラムの処理そのものを表す</li>
   <li>Data Gear はデータそのものを表す</li>
   <li>Code Gear は必要な Input Data Gear が揃ったら実行し、Output Data Gear を生成する</li>
@@ -156,55 +151,6 @@
 </div>
 <div class='slide '>
 <!-- _S9SLIDE_ -->
-<h2 id="continuation-based-c">Continuation based C</h2>
-<ul>
-  <li>Gears OS の実装は本研究室で開発している Continuation based C(CbC) を用いる</li>
-  <li>CbC は処理を Code Gear を用いて記述する事を基本とする</li>
-</ul>
-
-
-</div>
-<div class='slide '>
-<!-- _S9SLIDE_ -->
-<h2 id="continuation-based-c-1">Continuation based C</h2>
-<ul>
-  <li>Code Gear の定義は <code>__code CS名</code> で行う</li>
-  <li>Code Gear 間は <code>goto CS名</code> で移動する。この移動を継続と呼ぶ</li>
-  <li>Code Gear の継続は C の関数呼び出しとは異なり、戻り値を持たないためスタックの変更を行わない
-    <ul>
-      <li>このような環境を持たない継続を軽量継続と呼ぶ</li>
-    </ul>
-  </li>
-</ul>
-
-
-</div>
-<div class='slide '>
-<!-- _S9SLIDE_ -->
-<h2 id="continuation-based-c-2">Continuation based C</h2>
-<ul>
-  <li>このコードは cg0、cg1 の2つの Code Gear を定義している</li>
-  <li>cg0 内の <code>goto cg1</code> でgj1 への継続を行っている
-    <ul>
-      <li>ここで(a+b) が cg1 への入力になる</li>
-    </ul>
-  </li>
-</ul>
-
-<pre lang="c"><code>__code cg0(int a, int b) {
-    goto cg1(a+b);
-
-}
-
-__code cg1(int c) {
-    goto cg2(c);
-}
-</code></pre>
-
-
-</div>
-<div class='slide '>
-<!-- _S9SLIDE_ -->
 <h2 id="section-1">メタ計算</h2>
 <ul>
   <li>メタ計算 は通常の計算を実行するための計算</li>
@@ -219,7 +165,7 @@
 <!-- _S9SLIDE_ -->
 <h2 id="meta-gear">Meta Gear</h2>
 <ul>
-  <li>メタ計算 は Code Gearの接続の間に行われる</li>
+  <li>メタ計算 は Code Gear の接続間に行われる</li>
   <li>メタ計算 の処理も Code/Data Gear で実現する</li>
   <li>この Gear を Meta Code/Data Gearと呼ぶ
     <ul>
@@ -238,17 +184,53 @@
 </div>
 <div class='slide '>
 <!-- _S9SLIDE_ -->
+<h2 id="continuation-based-c">Continuation based C</h2>
+<ul>
+  <li>Gears OS の実装は本研究室で開発している Continuation based C(CbC) を用いる</li>
+  <li>CbC は Code Gear を用いて記述する事を基本とする</li>
+</ul>
+
+
+</div>
+<div class='slide '>
+<!-- _S9SLIDE_ -->
+<h2 id="continuation-based-c-1">Continuation based C</h2>
+<ul>
+  <li>CbC の Code Gear の定義は <strong>__code CG名</strong> で行う</li>
+  <li>Code Gear 間は <strong>goto CG名</strong> で移動する。この移動を継続と呼ぶ</li>
+  <li>Code Gear の継続は C の関数呼び出しとは異なり、戻り値を持たないためスタックの変更を行わない
+    <ul>
+      <li>このような環境を持たない継続を軽量継続と呼ぶ</li>
+    </ul>
+  </li>
+  <li>下記のコードでは Code Gear を2つ定義し、 cg0 から cg1 への継続を示している</li>
+</ul>
+
+<pre lang="c"><code>__code cg0(int a, int b) {
+    goto cg1(a+b);
+
+}
+
+__code cg1(int c) {
+    goto cg2(c);
+}
+</code></pre>
+
+
+</div>
+<div class='slide '>
+<!-- _S9SLIDE_ -->
 <h2 id="context">Context</h2>
 <ul>
   <li>Context は接続可能な Code/Data Gear の集合を表現する Meta Data Gear</li>
-  <li>従来のOS のスレッドやプロセスに対応
+  <li>従来のOS のスレッドやプロセスに対応し、以下の要素を定義している
     <ul>
-      <li>独立したメモリ空間を持つ</li>
+      <li>独立したメモリ空間</li>
       <li>Code Gear、 Data Gear へのポインタ</li>
       <li>並列実行用の Task 情報</li>
+      <li>Data Gear の型情報</li>
     </ul>
   </li>
-  <li>を持つ</li>
   <li>Gears OS ではメタ計算でこの Context を経由して Data Gear にアクセスする</li>
 </ul>
 
@@ -282,14 +264,27 @@
 </div>
 <div class='slide '>
 <!-- _S9SLIDE_ -->
-<h2 id="code-gear--stub-code-gear">Code Gear の stub Code Gear</h2>
+<h2 id="stub-code-gear">stub Code Gear</h2>
 <ul>
   <li>Data Gear にアクセスするにはContext を経由する</li>
   <li>だが、通常の Code Gear では Meta Data Gear である Context の参照は避ける必要がある</li>
   <li>Gears OS では通常の Code Gear で必要な Data Gear を Context から取り出す stub Code Gear を用意する</li>
 </ul>
 
-<pre lang="c"><code>
+<pre lang="c"><code>// normal level Code Gear
+__code cg0(struct Context* context, struct Integer integer, struct Queue queue) {
+    ...
+}
+
+// meta level stub Code Gear
+__code cg0_stub(struct Context* context) {
+    // get data index number
+    Integer integer = &amp;context-&gt;data[context-&gt;dataNum]-&gt;Integer
+    // get enum data
+    Queue* queue = &amp;context-&gt;data[Queue]-&gt;Queue;
+    // continuation Code Gear
+    goto cg0(context, integer, queue);
+}
 </code></pre>
 
 
@@ -298,9 +293,13 @@
 <!-- _S9SLIDE_ -->
 <h2 id="context--stub-code-gear-">Context での stub Code Gear の記述の問題点</h2>
 <ul>
-  <li>Context プログラム全体で使用する Code Gear と Data Gear の集合</li>
-  <li>全ての Code と Data の組合せをContext から 全て展開し、その組合せを stub Code Gear に書く必要がある</li>
-  <li>Gears OS を実装するに連れて、 stub Code Gear の記述が煩雑になる場所がでてきた</li>
+  <li>stub Code Gear は Context から Code Gear と Data Gear の全ての組合せを展開して記述する必要がある</li>
+  <li>しかし、 Gears OS を実装するに連れて、 stub Code Gear の記述が煩雑になる場所がでてきた
+    <ul>
+      <li>Data Gear は番号で指定するため、 Code Gear が どの Data Gear の番号に対応しているかを記述する必要がある</li>
+      <li>自動化するために、同じ番号の Data Gear  を使いまわす問題</li>
+    </ul>
+  </li>
   <li>そのため Gears OS のモジュール化する仕組みとして <strong>Interface</strong> を導入した</li>
 </ul>
 
@@ -311,8 +310,8 @@
 <h2 id="interface">Interface</h2>
 <ul>
   <li>Interface はある Data Gear と それに対する操作(API) を行う Code Gear の集合を表現する Meta Data Gear</li>
-  <li>stub Code Gear は実装した Code Gear で決まった形になるため、自動生成が可能である</li>
-  <li>Interface を導入することで、 Stack や Queue などのデータ構造を使用と実装に分けて記述することが出来る</li>
+  <li>stub Code Gear はInteface を実装した Code Gear で決まった形になるため、自動生成が可能である</li>
+  <li>Interface を導入することで、 Stack や Queue などのデータ構造を仕様と実装に分けて記述することが出来る</li>
   <li>Interface は Java のインターフェース、 Haskell の型クラスに対応する</li>
 </ul>
 
@@ -379,7 +378,7 @@
 <h2 id="interface--2">Interface の実装例</h2>
 <ul>
   <li>SingleLinkedQueue の put 実装</li>
-  <li>引数は Queue Interface の定義にあわせる</li>
+  <li>引数は Queue Interface の put 定義にあわせる</li>
   <li>第1引数は 実装対象の Data Gear の型になる</li>
   <li>第3引数の(…) は Output Data Gear を記述する
     <ul>
@@ -405,12 +404,7 @@
 <h2 id="interface--code-gear-">Interface を利用した Code Gear の呼び出し</h2>
 <ul>
   <li>Interface を利用した Code Gear への継続は <code>goto interface-&gt;method</code> で行われる</li>
-  <li>ここでの <strong>interface</strong>  は Interfaceの型で包んだポインタ、 <strong>method</strong> は実装した Code Gear に対応</li>
-  <li>この構文は実際にはスクリプトで変換される
-    <ul>
-      <li>変換後はメタレベルのコードになる</li>
-    </ul>
-  </li>
+  <li>ここでの <strong>interface</strong>  は Interfaceの型で包んだポインタ、 <strong>method</strong> は実装した Code Gear に対応する</li>
 </ul>
 
 <pre><code>__code code1() {
@@ -427,9 +421,14 @@
 <!-- _S9SLIDE_ -->
 <h2 id="interface--code-gear--1">Interface を利用した Code Gear の呼び出し(スクリプト変換後)</h2>
 <ul>
+  <li>Interface を利用した Code Gear の継続はスクリプトによって変換される
+    <ul>
+      <li>変換後は Context を参照するため、メタレベルの記述になる</li>
+    </ul>
+  </li>
   <li>Gearef マクロは Context から Interface の引数格納用の Data Gear を取り出す</li>
   <li>この Data Gear は Context を初期化した際に特別に生成され、型は Interface と同じになる</li>
-  <li>呼び出すCode Gear の引数情報に合わせて引数に格納</li>
+  <li>呼び出すCode Gear の引数情報に合わせて引数に格納し、 実装された Code Gear へ継続する</li>
 </ul>
 
 <pre><code>__code code1(struct Context *context) {
@@ -449,8 +448,8 @@
 <!-- _S9SLIDE_ -->
 <h2 id="interface--stub-code-gear">Interface での stub Code Gear</h2>
 <ul>
-  <li>メタ計算で格納された引数は、 stub Code Gear で Code Gear に渡される</li>
-  <li>Interface を実装した Code Gear は stub Code Gear の自動生成が可能である</li>
+  <li>メタ計算で格納された引数は stub Code Gear で Code Gear に渡される</li>
+  <li>Interface を実装した Code Gear は Interface の定義から stub Code Gear の自動生成が可能である</li>
 </ul>
 
 <pre lang="c"><code>__code putSingleLinkedQueue(struct Context *context,struct SingleLinkedQueue* queue, union Data* data, enum Code next) {
@@ -477,8 +476,8 @@
 <!-- _S9SLIDE_ -->
 <h2 id="section-2">並列処理の構成</h2>
 <ul>
-  <li>今回は並列処理を行う機構の実装を行う</li>
-  <li>構成要素
+  <li>今回はInterface を利用した並列処理機構の実装を行う</li>
+  <li>構成要素として以下が挙げられる
     <ul>
       <li>Task(Context)</li>
       <li>TaskManager
@@ -491,6 +490,11 @@
           <li>SynchronizedQueue から Task を取得し、実行する</li>
         </ul>
       </li>
+      <li>SynchronizedQueue
+        <ul>
+          <li>マルチスレッド 環境でもデータの一貫性を保証する Queue</li>
+        </ul>
+      </li>
     </ul>
   </li>
 </ul>
@@ -502,7 +506,7 @@
 <h2 id="taskcontext">Task(Context)</h2>
 <ul>
   <li>Gears OS では並列実行する Task を Context で表現する</li>
-  <li>Context Task用の情報として以下の情報をもっている
+  <li>Context は Task用の情報として以下の情報をもっている
     <ul>
       <li>実行する Code Gear</li>
       <li>Input/Output Data Gear の格納場所</li>
@@ -511,7 +515,7 @@
   </li>
   <li>実際に実行される Code Gear の引数情報は Interface の Code Gear 実装と同等に記述できる
     <ul>
-      <li>stub Code Gear は自動生成される</li>
+      <li>Code Gear の stub Code Gear は自動生成される</li>
     </ul>
   </li>
 </ul>
@@ -528,8 +532,8 @@
 <!-- _S9SLIDE_ -->
 <h2 id="taskmanger">TaskManger</h2>
 <ul>
-  <li>初期化時に決まった数の Worker を作成</li>
   <li>依存関係を解決した Task を各 Worker の Queue に送信する</li>
+  <li>Worker を作成、終了処理も行う</li>
 </ul>
 
 <div style="text-align: center;">
@@ -543,7 +547,7 @@
 <h2 id="worker">Worker</h2>
 <ul>
   <li>初期化の際にスレッドと Worker 用の Context を生成する</li>
-  <li>TaskManager から送信された Task を取得して実行する</li>
+  <li>生成されたスレッドではTaskManager から送信された Task を取得して実行する</li>
 </ul>
 
 <div>
@@ -560,6 +564,7 @@
             <li>Worker は再び Queue から Task を取得する</li>
         </ol>
     </div>
+    <div style="clear: both;"></div>
 </div>
 
 
@@ -569,9 +574,13 @@
 <h2 id="synchronized-queue">Synchronized Queue</h2>
 <ul>
   <li>Worker で使用される Queue</li>
-  <li>TaskManager を経由して Task を送信するスレッドと Task を取得するスレッドで操作される</li>
+  <li>Task を送信するスレッドと Task を取得するスレッドで操作される</li>
   <li>そのためマルチスレッドでのデータの一貫性を保証する必要がある</li>
-  <li>Gears OS では CAS(Check and Set、 Compare and Swap) を使用した Synchronized Queue として実装する</li>
+  <li>Gears OS では CAS(Check and Set、 Compare and Swap) を使用した Synchronized Queue として実装する
+    <ul>
+      <li>CAS は値を更新する際に更新前の値と実際に保存されているメモリ番地の値を比較する</li>
+    </ul>
+  </li>
   <li>この Queue は Queue Interface を実装し、 List を利用した実装を行った</li>
 </ul>
 
@@ -593,17 +602,30 @@
 <!-- _S9SLIDE_ -->
 <h2 id="section-3">依存関係の解決</h2>
 <ul>
-  <li>依存関係の解決は Data Gear にメタレベルで依存関係解決用の Queueをもたせることで行う</li>
-  <li>Code Gear を実行した後、 Output Data Gear を書き出す処理を行う</li>
-  <li>書き出し処理は Data Gear の Queue から依存関係にある Task を参照する</li>
-  <li>Task には実行に必要な Input Data Gear のカウンタを持っているため、そのカウンタをデクリメントする</li>
-  <li>カウンタが0になったら Input Data Gear が揃ったことになるため、TaskManager を経由して Worker に送信する</li>
+  <li>依存関係の解決は Data Gear がメタレベルで持っている Queue を使用する</li>
+  <li>この Queue には Data Gear に依存関係がある Code Gear が格納されている</li>
 </ul>
 
-<div style="text-align: center;">
-    <img src="./images/dependency.svg" alt="message" width="600" />
+<div>
+    <div style="float: left;">
+        <img src="./images/dependency.svg" alt="message" width="550" />
+    </div>
+    <div style="float: left; font-size=100%;">
+        <ol>
+            <li>Task の Code Gear を実行する</li>
+            <li>Output Data Gear の書き出し処理を行う</li>
+                この際にメタレベルの Queue を参照する
+            <li>依存関係にある Task を取り出し、 待っている</li>
+                Data Gearのカウンタをデクリメントする
+        </ol>
+    </div>
+    <div style="clear: both;"></div>
 </div>
 
+<ul>
+  <li>カウンタの値が0になった実行可能な Task は TaskManager を通して Worker に送信される</li>
+</ul>
+
 
 </div>
 <div class='slide '>
@@ -612,7 +634,7 @@
 <ul>
   <li>並列実行の Task の生成は新しく Context を生成し、実行する Code Gear、待ち合わせる Input Data Gear の数、Input/Output Data Gear への参照を設定する</li>
   <li>この記述を直接書くと Meta Data Gear である Context を直接参照しているため、ノーマルレベルでの記述は好ましくない</li>
-  <li>Task の設定の記述は煩雑な記述であるが、並列実行サれることを除けば通常の CbC の goto 文と同等である</li>
+  <li>Task の設定の記述は煩雑な記述であるが、並列実行されることを除けば通常の CbC の goto 文と同等である</li>
   <li>そこで Context を直接参照しない並列構文、 <strong>par goto</strong> 構文を新たに考案した</li>
   <li>par goto 構文には引数として Input/Output Data Gear等を渡す
     <ul>
@@ -641,9 +663,15 @@
   <li>今回 CUDAWorker、CUDAExecutor、 CUDABuffer を使用して CUDA に合わせた Code Gear を提供する</li>
 </ul>
 
-<div style="text-align: center;">
-    <img src="./images/cudaArchitecture.svg" alt="message" width="500" />
+
 </div>
+<div class='slide '>
+<!-- _S9SLIDE_ -->
+<h2 id="cudaworker">CUDAWorker</h2>
+<ul>
+  <li>CUDA で実行する Task を受け取る Worker</li>
+  <li>初期化の際に CUDA ライブラリの初期化等を行う</li>
+</ul>
 
 
 </div>
@@ -651,7 +679,7 @@
 <!-- _S9SLIDE_ -->
 <h2 id="cudaexecutor">CUDAExecutor</h2>
 <ul>
-  <li>CUDA Executor は Executor Interface を実装した以下の Code Gear を持つ
+  <li>CUDAExecutor は Executor Interface を実装した以下の Code Gear を持つ
     <ul>
       <li>HostからDevice へのデータの送信(read)</li>
       <li>kernel の実行(exec)</li>
@@ -677,7 +705,7 @@
 <!-- _S9SLIDE_ -->
 <h2 id="cudabuffer">CUDABuffer</h2>
 <ul>
-  <li>Host、 Device 間でデータのやり取りをする際、 Gears OS での Data Gear をDevice 用にマッピングする必要がある
+  <li>Host、Device 間でデータのやり取りをする際、 Gears OS での Data Gear をDevice 用にマッピングする必要がある
     <ul>
       <li>Device にデータ領域を確保するにはサイズの指定が必要</li>
       <li>Data Gear には Meta Data Gear でデータのサイズを持っている</li>
@@ -739,9 +767,6 @@
 <ul>
   <li>Twice は与えられた整数配列を2倍にする例題である</li>
   <li>並列実行の依存関係がなく、並列度が高い課題である</li>
-  <li>要素数 2^27</li>
-  <li>CPU での実行時は 2^27 を 2^6 個に分割して Task を生成する</li>
-  <li>GPU での実行時は1次元の block 数を 2^15、 block 内の thread 数を 2^10 で展開</li>
 </ul>
 
 
@@ -749,6 +774,15 @@
 <div class='slide '>
 <!-- _S9SLIDE_ -->
 <h2 id="twice-">Twice の結果</h2>
+<ul>
+  <li>要素数 2^27</li>
+  <li>CPU での実行時は 2^27 を 2^6 個に分割して Task を生成する</li>
+  <li>GPU での実行時は1次元の block 数を 2^15、 block 内の thread 数を 2^10 で展開</li>
+  <li>1 CPU と 32 CPU では 約27.1倍の速度向上が見られた</li>
+  <li>GPU 実行は kernel のみの実行時間は32 CPU に比べて約7.2倍の速度向上、通信時間を含めると 16 CPU より遅い</li>
+  <li>通信時間がオーバーヘッドになっている</li>
+</ul>
+
 <table border="1" align="center" width="50%">
   <tbody>
     <tr>
@@ -781,7 +815,7 @@
     </tr>
     <tr>
       <td style="text-align: center;">GPU</td>
-      <td style="text-align: right;">43.496</td>
+      <td style="text-align: right;">127.018</td>
     </tr>
     <tr>
       <td style="text-align: center;">GPU(kernel only)</td>
@@ -790,16 +824,6 @@
   </tbody>
 </table>
 
-<ul>
-  <li>1 CPU と 32 CPU では 約27.1倍の速度向上が見られた</li>
-  <li>GPU での実行は kernel のみの実行時間は32 CPU に比べて約7.2倍の速度向上
-    <ul>
-      <li>通信時間を含めると 16 CPU より遅い</li>
-      <li>通信時間がオーバーヘッドになっている</li>
-    </ul>
-  </li>
-</ul>
-
 
 </div>
 <div class='slide '>
@@ -808,9 +832,6 @@
 <ul>
   <li>並列処理向けのソートアルゴリズム</li>
   <li>決まった2点間の要素の入れ替えをステージ毎に並列に実行し、 Output Data Gear として書き出し、次のステージの Code Gear の Input Data Gear とする</li>
-  <li>要素数 2^24</li>
-  <li>CPU での実行時は 2^24 を 2^6 個に分割して Task を生成する</li>
-  <li>GPU での実行時は1次元の block 数を 2^14、 block 内の thread 数を 2^10 で展開</li>
 </ul>
 
 <div style="text-align: center;">
@@ -822,6 +843,14 @@
 <div class='slide '>
 <!-- _S9SLIDE_ -->
 <h2 id="bitonicsort-">BitonicSort の結果</h2>
+<ul>
+  <li>要素数 2^24</li>
+  <li>CPU での実行時は 2^24 を 2^6 個に分割して Task を生成する</li>
+  <li>GPU での実行時は1次元の block 数を 2^14、 block 内の thread 数を 2^10 で展開</li>
+  <li>1 CPU と 32 CPU で約22.12倍の速度向上</li>
+  <li>GPU は通信時間を含めると 8 CPU の約1.16倍、 kernel のみの実行では 32 CPU の約11.48倍になった</li>
+  <li>現在の Gears OS の CUDA 実装では Output Data Gear を書き出す際に一度 GPU から CPU へ kernel の結果の書き出しを行っているため、差がでてしまった</li>
+</ul>
 
 <table border="1" align="center" width="50%">
   <tbody>
@@ -864,18 +893,81 @@
   </tbody>
 </table>
 
+
+</div>
+<div class='slide '>
+<!-- _S9SLIDE_ -->
+<h2 id="openmp-">OpenMP との比較</h2>
 <ul>
-  <li>1 CPU と 32 CPU で約22.12倍の速度向上</li>
-  <li>GPU は通信時間を含めると 8 CPU の役1.16倍、 kernel のみの実行では 32 CPU の役11.48倍になった</li>
-  <li>現在の Gears OS の CUDA 実装では Output Data Gear を書き出す際に一度 GPU から CPU へ kernel の結果の書き出しを行っているため、差がでてしまった</li>
+  <li>OpenMP は C、 C++ のプログラムにアノテーションを付けることで並列化を行う</li>
+  <li>データの待ち合わせ処理はバリア等のアノテーションで記述する</li>
+  <li>Gears OS は並列処理を par goto 構文、 データの待ち合わせを Code Gear と Input/Ouput Data Gear の関係で行う</li>
 </ul>
 
+<pre lang="c"><code>#pragma omp parallel for
+for(int i = 0; i &lt; length; i++) {
+    a[i] = a[i] * 2;
+}
+</code></pre>
+
+
+</div>
+<div class='slide '>
+<!-- _S9SLIDE_ -->
+<h2 id="openmp--1">OpenMP との比較</h2>
+<ul>
+  <li>OpenMP で Twice を実装し、速度比較を行った</li>
+  <li>OpenMP は 1CPU と 32CPU で約10.8倍の速度向上が見られた</li>
+  <li>一方 Gears OS では約27.1倍と台数効果は高くなっている</li>
+  <li>しかし、 Gears OS は 1CPU の実行速度が OpenMP に比べて大幅に遅くなっている</li>
+</ul>
+
+<div style="text-align: center;">
+    <img src="./images/vsopenmp.svg" alt="message" width="500" />
+</div>
+
 
 </div>
 <div class='slide '>
 <!-- _S9SLIDE_ -->
-<h2 id="openmp-">OpenMP との比較</h2>
-<p>## Go との比較</p>
+<h2 id="go-">Go 言語との比較</h2>
+<ul>
+  <li>Go 言語は並列実行を <strong>go funciton(argv)</strong> の構文で行う。 この実行を goroutine と呼ぶ</li>
+  <li>データの待ち合わせはチャネルというデータ構造で行う</li>
+  <li>チャネルでのデータの送受信は <strong>&lt;-</strong> を使用して行うため、簡潔に書くことが出来る</li>
+  <li>しかし、 チャネルは複数の goroutine で共有されるため、データの送信元が推測しづらい</li>
+  <li>Gears OS では goroutine は par goto 文とほぼ同等に扱える</li>
+  <li>par goto 文では書き出す Data Gear を指定するため、書き出し元が推測しやすい</li>
+</ul>
+
+<pre lang="go"><code>c := make(chan []int)
+for i :=0; i &lt; *split; i++ {
+    // call goroutine
+    go twice(list, prefix, i, c);
+}
+
+func twice(list []int, prefix int, index int, c chan []int) {
+    for i := 0; i &lt; prefix; i++ {
+        list[prefix*index+i] = list[prefix*index+i] * 2;
+    }
+    c &lt;- list
+}
+</code></pre>
+
+
+</div>
+<div class='slide '>
+<!-- _S9SLIDE_ -->
+<h2 id="go--1">Go 言語との比較</h2>
+<ul>
+  <li>Go 言語でも OpenMP と同様に Twice を実装し、速度比較を行った</li>
+  <li>Go 言語は 1CPU と 32CPU で約4.33倍の速度向上が見られた</li>
+  <li>OpenMP と同様に台数効果自体は Gears OS が高いが、 1CPU での実行時間は Go 言語が大幅に速い</li>
+</ul>
+
+<div style="text-align: center;">
+    <img src="./images/vsgo.svg" alt="message" width="500" />
+</div>
 
 
 </div>
@@ -898,7 +990,7 @@
   <li>Gears OS の並列処理の信頼性の保証、チューニングを行う</li>
   <li>Gears OS では検証とモデル検査をメタレベルで実現することで信頼性を保証する
     <ul>
-      <li>証明はCbC のプログラムヲ証明支援系の Agda に対応して行う。 並列処理の信頼性を保証するには SynchronizedQueue の証明を行う必要がある</li>
+      <li>証明は CbC のプログラムを証明支援系の Agda に対応して行う。 並列処理の信頼性を保証するには SynchronizedQueue の証明を行う必要がある</li>
       <li>モデル検査は CbC で記述された モデル検査器である akasha を使用して行う。 モデル検査の方針としては Code Gear の並列実行を擬似並列で実行し、全ての組合せを列挙する方法で行う</li>
     </ul>
   </li>
--- a/slide/slide.md	Mon Feb 12 03:16:15 2018 +0900
+++ b/slide/slide.md	Mon Feb 12 12:07:12 2018 +0900
@@ -56,11 +56,11 @@
 - CbC は Code Gear を用いて記述する事を基本とする
 
 ## Continuation based C
-- Code Gear の定義は ``__code CS名`` で行う
-- Code Gear 間は ``goto CS名`` で移動する。この移動を継続と呼ぶ
+- CbC の Code Gear の定義は **__code CG名** で行う
+- Code Gear 間は **goto CG名** で移動する。この移動を継続と呼ぶ
 - Code Gear の継続は C の関数呼び出しとは異なり、戻り値を持たないためスタックの変更を行わない
     - このような環境を持たない継続を軽量継続と呼ぶ
-- 下のコード例では Code Gear を2つ定義し、 cg0 から cg1 への継続を示している
+- 下記のコードでは Code Gear を2つ定義し、 cg0 から cg1 への継続を示している
 
 ``` c
 __code cg0(int a, int b) {
@@ -103,7 +103,7 @@
 };
 ```
 
-## Code Gear の stub Code Gear
+## stub Code Gear
 - Data Gear にアクセスするにはContext を経由する
 - だが、通常の Code Gear では Meta Data Gear である Context の参照は避ける必要がある
 - Gears OS では通常の Code Gear で必要な Data Gear を Context から取り出す stub Code Gear を用意する
@@ -126,15 +126,15 @@
 ```
 
 ## Context での stub Code Gear の記述の問題点
-- stub Code Gear は Context から Code Gear と Data Gear の組合せを展開して記述する必要がある
+- stub Code Gear は Context から Code Gear と Data Gear の全ての組合せを展開して記述する必要がある
 - しかし、 Gears OS を実装するに連れて、 stub Code Gear の記述が煩雑になる場所がでてきた
-    - Context は全ての Code と Data の集合を表現しているため、全て組合せを展開する必要がある
-    - 
+    - Data Gear は番号で指定するため、 Code Gear が どの Data Gear の番号に対応しているかを記述する必要がある
+    - 自動化するために、同じ番号の Data Gear  を使いまわす問題
 - そのため Gears OS のモジュール化する仕組みとして **Interface** を導入した
 
 ## Interface
 - Interface はある Data Gear と それに対する操作(API) を行う Code Gear の集合を表現する Meta Data Gear
-- stub Code Gear は実装した Code Gear で決まった形になるため、自動生成が可能である
+- stub Code Gear はInteface を実装した Code Gear で決まった形になるため、自動生成が可能である
 - Interface を導入することで、 Stack や Queue などのデータ構造を仕様と実装に分けて記述することが出来る
 - Interface は Java のインターフェース、 Haskell の型クラスに対応する
 
@@ -183,7 +183,7 @@
 
 ## Interface の実装例
 - SingleLinkedQueue の put 実装
-- 引数は Queue Interface の定義にあわせる
+- 引数は Queue Interface の put 定義にあわせる
 - 第1引数は 実装対象の Data Gear の型になる
 - 第3引数の(...) は Output Data Gear を記述する
     -  ... は可変長引数のような扱いで、 継続先の Code Gear が複数の値をInput Data Gear とする可能性がある
@@ -201,9 +201,7 @@
 
 ## Interface を利用した Code Gear の呼び出し
 - Interface を利用した Code Gear への継続は `goto interface->method` で行われる
-- ここでの **interface**  は Interfaceの型で包んだポインタ、 **method** は実装した Code Gear に対応
-- この構文は実際にはスクリプトで変換される
-    - 変換後はメタレベルのコードになる
+- ここでの **interface**  は Interfaceの型で包んだポインタ、 **method** は実装した Code Gear に対応する
 
 ```
 __code code1() {
@@ -215,9 +213,11 @@
 ```
 
 ## Interface を利用した Code Gear の呼び出し(スクリプト変換後)
+- Interface を利用した Code Gear の継続はスクリプトによって変換される
+    - 変換後は Context を参照するため、メタレベルの記述になる
 - Gearef マクロは Context から Interface の引数格納用の Data Gear を取り出す
 - この Data Gear は Context を初期化した際に特別に生成され、型は Interface と同じになる
-- 呼び出すCode Gear の引数情報に合わせて引数に格納
+- 呼び出すCode Gear の引数情報に合わせて引数に格納し、 実装された Code Gear へ継続する
 
 ```
 __code code1(struct Context *context) {
@@ -232,8 +232,8 @@
 ```
 
 ## Interface での stub Code Gear
-- メタ計算で格納された引数は、 stub Code Gear で Code Gear に渡される
-- Interface を実装した Code Gear は stub Code Gear の自動生成が可能である
+- メタ計算で格納された引数は stub Code Gear で Code Gear に渡される
+- Interface を実装した Code Gear は Interface の定義から stub Code Gear の自動生成が可能である
 
 ``` c
 __code putSingleLinkedQueue(struct Context *context,struct SingleLinkedQueue* queue, union Data* data, enum Code next) {
@@ -255,22 +255,24 @@
 ```
 
 ## 並列処理の構成
-- 今回は並列処理を行う機構の実装を行う
-- 構成要素
+- 今回はInterface を利用した並列処理機構の実装を行う
+- 構成要素として以下が挙げられる
     - Task(Context)
     - TaskManager
         - Worker の生成、依存関係を解決したTask を Worker に送信する
     - Worker
         - SynchronizedQueue から Task を取得し、実行する
+    - SynchronizedQueue
+        - マルチスレッド 環境でもデータの一貫性を保証する Queue
 
 ## Task(Context)
 - Gears OS では並列実行する Task を Context で表現する 
-- Context Task用の情報として以下の情報をもっている
+- Context は Task用の情報として以下の情報をもっている
     - 実行する Code Gear
     - Input/Output Data Gear の格納場所
     - 待っている Input Data Gear の数
 - 実際に実行される Code Gear の引数情報は Interface の Code Gear 実装と同等に記述できる
-    - stub Code Gear は自動生成される
+    - Code Gear の stub Code Gear は自動生成される
 
 ``` c
 __code add(struct Integer* input1, struct Integer* input2, __code next(struct Integer* output, ...)) {
@@ -280,8 +282,8 @@
 ```
 
 ## TaskManger
-- 初期化時に決まった数の Worker を作成
 - 依存関係を解決した Task を各 Worker の Queue に送信する
+- Worker を作成、終了処理も行う
 
 <div style="text-align: center;">
     <img src="./images/sendTask.svg" alt="message" width="800">
@@ -289,7 +291,7 @@
 
 ## Worker
 - 初期化の際にスレッドと Worker 用の Context を生成する 
-- TaskManager から送信された Task を取得して実行する
+- 生成されたスレッドではTaskManager から送信された Task を取得して実行する
 
 <div>
     <div style="float: left;">
@@ -305,13 +307,15 @@
             <li>Worker は再び Queue から Task を取得する</li>
         </ol>
     </div>
+    <div style="clear: both;"></div>
 </div>
 
 ## Synchronized Queue
 - Worker で使用される Queue
-- TaskManager を経由して Task を送信するスレッドと Task を取得するスレッドで操作される
+- Task を送信するスレッドと Task を取得するスレッドで操作される
 - そのためマルチスレッドでのデータの一貫性を保証する必要がある
 - Gears OS では CAS(Check and Set、 Compare and Swap) を使用した Synchronized Queue として実装する
+    - CAS は値を更新する際に更新前の値と実際に保存されているメモリ番地の値を比較する
 - この Queue は Queue Interface を実装し、 List を利用した実装を行った
 
 ```
@@ -328,16 +332,27 @@
 ```
 
 ## 依存関係の解決
-- 依存関係の解決は Data Gear にメタレベルで依存関係解決用の Queueをもたせることで行う
-- Code Gear を実行した後、 Output Data Gear を書き出す処理を行う
-- 書き出し処理は Data Gear の Queue から依存関係にある Task を参照する
-- Task には実行に必要な Input Data Gear のカウンタを持っているため、そのカウンタをデクリメントする
-- カウンタが0になったら Input Data Gear が揃ったことになるため、TaskManager を経由して Worker に送信する
+- 依存関係の解決は Data Gear がメタレベルで持っている Queue を使用する
+- この Queue には Data Gear に依存関係がある Code Gear が格納されている
 
-<div style="text-align: center;">
-    <img src="./images/dependency.svg" alt="message" width="600">
+<div>
+    <div style="float: left;">
+        <img src="./images/dependency.svg" alt="message" width="550">
+    </div>
+    <div style="float: left; font-size=100%;">
+        <ol>
+            <li>Task の Code Gear を実行する</li>
+            <li>Output Data Gear の書き出し処理を行う</li>
+                この際にメタレベルの Queue を参照する
+            <li>依存関係にある Task を取り出し、 待っている</li>
+                Data Gearのカウンタをデクリメントする
+        </ol>
+    </div>
+    <div style="clear: both;"></div>
 </div>
 
+- カウンタの値が0になった実行可能な Task は TaskManager を通して Worker に送信される
+
 ## 並列構文
 - 並列実行の Task の生成は新しく Context を生成し、実行する Code Gear、待ち合わせる Input Data Gear の数、Input/Output Data Gear への参照を設定する
 - この記述を直接書くと Meta Data Gear である Context を直接参照しているため、ノーマルレベルでの記述は好ましくない
@@ -361,12 +376,12 @@
 - CUDA は処理の最小の単位を thread とし、それをまとめた block を展開し Device 上で実行されるプログラム(Kernel)を実行する
 - 今回 CUDAWorker、CUDAExecutor、 CUDABuffer を使用して CUDA に合わせた Code Gear を提供する
 
-<div style="text-align: center;">
-    <img src="./images/cudaArchitecture.svg" alt="message" width="500">
-</div>
+## CUDAWorker
+- CUDA で実行する Task を受け取る Worker
+- 初期化の際に CUDA ライブラリの初期化等を行う
 
 ## CUDAExecutor
-- CUDA Executor は Executor Interface を実装した以下の Code Gear を持つ
+- CUDAExecutor は Executor Interface を実装した以下の Code Gear を持つ
     - HostからDevice へのデータの送信(read)
     - kernel の実行(exec)
     - Device から Host へのデータの書き出し(write)
@@ -384,7 +399,7 @@
 ```
 
 ## CUDABuffer
-- Host、 Device 間でデータのやり取りをする際、 Gears OS での Data Gear をDevice 用にマッピングする必要がある
+- Host、Device 間でデータのやり取りをする際、 Gears OS での Data Gear をDevice 用にマッピングする必要がある
     - Device にデータ領域を確保するにはサイズの指定が必要
     - Data Gear には Meta Data Gear でデータのサイズを持っている
     - だが、Data Gear の要素の中に Data Gear へのポインタがあるとポインタ分でサイズ計算してしまうため、 GPU では参照できなくなってしまう
@@ -416,11 +431,15 @@
 ## Twice
 - Twice は与えられた整数配列を2倍にする例題である
 - 並列実行の依存関係がなく、並列度が高い課題である
+
+## Twice の結果
 - 要素数 2^27
 - CPU での実行時は 2^27 を 2^6 個に分割して Task を生成する
 - GPU での実行時は1次元の block 数を 2^15、 block 内の thread 数を 2^10 で展開
+- 1 CPU と 32 CPU では 約27.1倍の速度向上が見られた
+- GPU 実行は kernel のみの実行時間は32 CPU に比べて約7.2倍の速度向上、通信時間を含めると 16 CPU より遅い
+- 通信時間がオーバーヘッドになっている
 
-## Twice の結果
 <table  border="1" align='center' width='50%'>
   <tbody>
     <tr>
@@ -462,22 +481,21 @@
   </tbody>
 </table>
 
-- 1 CPU と 32 CPU では 約27.1倍の速度向上が見られた
-- GPU 実行は kernel のみの実行時間は32 CPU に比べて約7.2倍の速度向上、通信時間を含めると 16 CPU より遅い
-- 通信時間がオーバーヘッドになっている
-
 ## BitonicSort
 - 並列処理向けのソートアルゴリズム
 - 決まった2点間の要素の入れ替えをステージ毎に並列に実行し、 Output Data Gear として書き出し、次のステージの Code Gear の Input Data Gear とする
-- 要素数 2^24
-- CPU での実行時は 2^24 を 2^6 個に分割して Task を生成する
-- GPU での実行時は1次元の block 数を 2^14、 block 内の thread 数を 2^10 で展開
 
 <div style="text-align: center;">
     <img src="./images/bitonicNetwork.svg" alt="message" width="500">
 </div>
 
 ## BitonicSort の結果
+- 要素数 2^24
+- CPU での実行時は 2^24 を 2^6 個に分割して Task を生成する
+- GPU での実行時は1次元の block 数を 2^14、 block 内の thread 数を 2^10 で展開
+- 1 CPU と 32 CPU で約22.12倍の速度向上
+- GPU は通信時間を含めると 8 CPU の約1.16倍、 kernel のみの実行では 32 CPU の約11.48倍になった
+- 現在の Gears OS の CUDA 実装では Output Data Gear を書き出す際に一度 GPU から CPU へ kernel の結果の書き出しを行っているため、差がでてしまった
 
 <table  border="1" align='center' width='50%'>
   <tbody>
@@ -520,12 +538,60 @@
   </tbody>
 </table>
 
-- 1 CPU と 32 CPU で約22.12倍の速度向上
-- GPU は通信時間を含めると 8 CPU の約1.16倍、 kernel のみの実行では 32 CPU の約11.48倍になった
-- 現在の Gears OS の CUDA 実装では Output Data Gear を書き出す際に一度 GPU から CPU へ kernel の結果の書き出しを行っているため、差がでてしまった
+
+## OpenMP との比較
+- OpenMP は C、 C++ のプログラムにアノテーションを付けることで並列化を行う
+- データの待ち合わせ処理はバリア等のアノテーションで記述する
+- Gears OS は並列処理を par goto 構文、 データの待ち合わせを Code Gear と Input/Ouput Data Gear の関係で行う
+
+``` c
+#pragma omp parallel for
+for(int i = 0; i < length; i++) {
+    a[i] = a[i] * 2;
+}
+```
 
 ## OpenMP との比較
-## Go との比較
+- OpenMP で Twice を実装し、速度比較を行った
+- OpenMP は 1CPU と 32CPU で約10.8倍の速度向上が見られた
+- 一方 Gears OS では約27.1倍と台数効果は高くなっている
+- しかし、 Gears OS は 1CPU の実行速度が OpenMP に比べて大幅に遅くなっている
+
+<div style="text-align: center;">
+    <img src="./images/vsopenmp.svg" alt="message" width="500">
+</div>
+
+## Go 言語との比較
+- Go 言語は並列実行を **go funciton(argv)** の構文で行う。 この実行を goroutine と呼ぶ
+- データの待ち合わせはチャネルというデータ構造で行う
+- チャネルでのデータの送受信は **<-** を使用して行うため、簡潔に書くことが出来る
+- しかし、 チャネルは複数の goroutine で共有されるため、データの送信元が推測しづらい
+- Gears OS では goroutine は par goto 文とほぼ同等に扱える
+- par goto 文では書き出す Data Gear を指定するため、書き出し元が推測しやすい
+
+``` go
+c := make(chan []int)
+for i :=0; i < *split; i++ {
+    // call goroutine
+    go twice(list, prefix, i, c);
+}
+
+func twice(list []int, prefix int, index int, c chan []int) {
+    for i := 0; i < prefix; i++ {
+        list[prefix*index+i] = list[prefix*index+i] * 2;
+    }
+    c <- list
+}
+```
+
+## Go 言語との比較
+- Go 言語でも OpenMP と同様に Twice を実装し、速度比較を行った
+- Go 言語は 1CPU と 32CPU で約4.33倍の速度向上が見られた
+- OpenMP と同様に台数効果自体は Gears OS が高いが、 1CPU での実行時間は Go 言語が大幅に速い
+
+<div style="text-align: center;">
+    <img src="./images/vsgo.svg" alt="message" width="500">
+</div>
 
 ## まとめ
 - Gears OS の並列処理機構の実装を行った
@@ -536,7 +602,7 @@
 ## 今後の課題
 - Gears OS の並列処理の信頼性の保証、チューニングを行う
 - Gears OS では検証とモデル検査をメタレベルで実現することで信頼性を保証する
-    - 証明はCbC のプログラムヲ証明支援系の Agda に対応して行う。 並列処理の信頼性を保証するには SynchronizedQueue の証明を行う必要がある
+    - 証明は CbC のプログラムを証明支援系の Agda に対応して行う。 並列処理の信頼性を保証するには SynchronizedQueue の証明を行う必要がある
     - モデル検査は CbC で記述された モデル検査器である akasha を使用して行う。 モデル検査の方針としては Code Gear の並列実行を擬似並列で実行し、全ての組合せを列挙する方法で行う
 - OpenMP、 Goとの比較から、 Gears OS が 1CPU での動作が遅いということがわかった。 
     - par goto 文を使用する度に Context を生成するため、 ある程度の時間がかかってしまう