Mercurial > hg > Papers > 2018 > parusu-master

view slide/slide.html @ 103:05d7669c4a29

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Fix
author Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date Tue, 13 Feb 2018 13:22:51 +0900
parents c48b78793176
children 4b49908418e2
line wrap: on
line source

<!DOCTYPE html>
<html>
<head>
   <meta http-equiv="content-type" content="text/html;charset=utf-8">
   <title>Gears OS の並列処理</title>

<meta name="generator" content="Slide Show (S9) v2.5.0 on Ruby 2.3.0 (2015-12-25) [x86_64-darwin16]">
<meta name="author"    content="伊波 立樹" >

<!-- style sheet links -->
<link rel="stylesheet" href="s6/themes/projection.css"   media="screen,projection">
<link rel="stylesheet" href="s6/themes/screen.css"       media="screen">
<link rel="stylesheet" href="s6/themes/print.css"        media="print">
<link rel="stylesheet" href="s6/themes/blank.css"        media="screen,projection">

<!-- JS -->
<script src="s6/js/jquery-1.11.3.min.js"></script>
<script src="s6/js/jquery.slideshow.js"></script>
<script src="s6/js/jquery.slideshow.counter.js"></script>
<script src="s6/js/jquery.slideshow.controls.js"></script>
<script src="s6/js/jquery.slideshow.footer.js"></script>
<script src="s6/js/jquery.slideshow.autoplay.js"></script>

<!-- prettify -->
<link rel="stylesheet" href="scripts/prettify.css">
<script src="scripts/prettify.js"></script>

<script>
  $(document).ready( function() {
    Slideshow.init();

    $('code').each(function(_, el) {
      if (!el.classList.contains('noprettyprint')) {
        el.classList.add('prettyprint');
        el.style.display = 'block';
      }
    });
    prettyPrint();
  } );

  
</script>

<!-- Better Browser Banner for Microsoft Internet Explorer (IE) -->
<!--[if IE]>
<script src="s6/js/jquery.microsoft.js"></script>
<![endif]-->



</head>
<body>

<div class="layout">
  <div id="header"></div>
  <div id="footer">
    <div align="right">
      <img src="s6/images/logo.svg" width="200px">
    </div>
  </div>
</div>

<div class="presentation">

  <div class='slide cover'>
    <table width="90%" height="90%" border="0" align="center">
      <tr>
        <td>
          <div align="center">
            <h1><font color="#808db5">Gears OS の並列処理</font></h1>
          </div>
        </td>
      </tr>
      <tr>
        <td>
          <div align="left">
            伊波 立樹
            琉球大学理工学研究科 河野研
            <hr style="color:#ffcc00;background-color:#ffcc00;text-align:left;border:none;width:100%;height:0.2em;">
          </div>
        </td>
      </tr>
    </table>
  </div>

<div class='slide '>
<!-- === begin markdown block ===

      generated by markdown/1.2.0 on Ruby 2.3.0 (2015-12-25) [x86_64-darwin16]
                on 2018-02-13 12:53:11 +0900 with Markdown engine kramdown (1.13.2)
                  using options {}
  -->

<!-- _S9SLIDE_ -->
<h2 id="section">並列処理の重要性</h2>
<ul>
  <li>並列処理は現在主流のマルチコアCPU の性能を発揮するには重要なものになっている</li>
  <li>しかし、並列処理のチューニングや信頼性を保証するのは難しい
    <ul>
      <li>共通資源の競合などの非決定的な実行が発生するため、従来のテストやデバッグではテストしきれない部分が残ってしまう</li>
      <li>GPU などのアーキテクチャに合わせた並列プログラミングの記述</li>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="gears-os">Gears OS</h2>
<ul>
  <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>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="gears-os-1">Gears OS</h2>
<ul>
  <li>本研究では Gears OS の並列処理機構、並列処理構文(par goto)の実装、Gears OS を実装するにつれて必要なったモジュール化の導入を行う</li>
  <li>また、並列処理を行う例題を用いて評価、 OpenMP、 Go 言語との比較を行う</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="code-geardata-gear">Code Gear/Data Gear</h2>
<ul>
  <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>
  <li>Code Gear と Input/Output Data Gear の対応から依存関係を解決し、Input Data Gear が揃った Code Gear の並列実行を行う</li>
</ul>

<div style="text-align: center;">
    <img src="./images/codegear-datagear-dependency.svg" alt="message" width="600" />
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-1">メタ計算</h2>
<ul>
  <li>メタ計算 は通常の計算を実行するための計算</li>
  <li>信頼性の確保やメモリ管理、スレッド管理、CPU、GPU の資源管理等</li>
  <li>Gears OS のメタ計算は通常の計算とは別の階層のメタレベルで行われる</li>
  <li>メタレベルは Code/Data Gear に対応して Meta Code/Data Gear で行われる</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="meta-gear">Meta Gear</h2>
<ul>
  <li>メタ計算 は Code Gear の接続間に行われる</li>
  <li>この Gear を Meta Code/Data Gearと呼ぶ
    <ul>
      <li>Meta Code Gear は メタ計算 のプログラム部分</li>
      <li>Meta Data Gear は Meta Code Gear で管理されるデータ部分</li>
    </ul>
  </li>
  <li>Gears OS は通常の Code/Data Gear から Meta Code/Data Gear 部分は見えないように実装を行う</li>
</ul>

<div style="text-align: center;">
    <img src="./images/meta_cg_dg.svg" alt="message" width="850" />
</div>


</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="data-gear-">Data Gear の表現</h2>
<ul>
  <li>Data Gear は構造体を用いて定義する</li>
  <li>メタ計算では任意の Data Gear を一律に扱うため、全ての Data Gear は共用体の中で定義される</li>
  <li>Data Gear をメモリに確保する際のサイズ情報はこの型から決定する</li>
</ul>

<pre lang="c"><code>/* data Gear define */
union Data {
    struct Timer {
        union Data* timer;
        enum Code start;
        enum Code end;
        enum Code next;
    } Timer;
    struct TimerImpl {
        double time;
    } TimerImpl;
    ....
};
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="context">Context</h2>
<ul>
  <li>Context は従来のOS のスレッドやプロセスに対応し、以下の要素をもっている Meta Data Gear
    <ul>
      <li>Data Gear を確保するためのメモリ空間</li>
      <li>Code Gear の名前と関数ポインタとの対応表
        <ul>
          <li>Code Gear は番号(enum)で指定する</li>
        </ul>
      </li>
      <li>Code Gear が参照する Data Gear へのポインタ
        <ul>
          <li>Code Gear と同じく Data Gear も番号で指定する</li>
        </ul>
      </li>
      <li>並列実行用の Task 情報</li>
      <li>Data Gear の型情報</li>
    </ul>
  </li>
  <li>Gears OS ではメタ計算で Context を経由して Code/Data Gear に番号でアクセスする</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<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>

<div style="text-align: center;">
    <img src="./images/contextContinuation.svg" alt="message" width="600" />
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="interface">Interface</h2>
<ul>
  <li>Gears OS を実装するに連れて、stub Code Gear の記述が煩雑になることがわかった</li>
  <li>そこで、Gears OS のモジュール化する仕組みとして <strong>Interface</strong> を導入した</li>
  <li>Interface はある Data Gear と それに対する操作(API) を行う Code Gear の集合を表現する Meta Data Gear</li>
  <li>stub Code Gear はInteface を実装した Code Gear で決まった形になるため、自動生成が可能</li>
  <li>Interface を導入することで、 Stack や Queue などのデータ構造を仕様と実装に分けて記述することが出来る</li>
  <li>Interface は Java のインターフェース、 Haskell の型クラスに対応する</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="interface-">Interface の定義</h2>
<ul>
  <li>Interface の定義には以下の内容を定義する
    <ul>
      <li>操作(API)の引数群の型</li>
      <li>操作(API)自体のCode Gear の型</li>
    </ul>
  </li>
</ul>

<pre lang="c"><code>typedef struct Queue&lt;Impl&gt;{
        // Data Gear parameter
        union Data* queue;
        union Data* data;
        __code next(...);
        __code whenEmpty(...);

        // Code Gear
        __code clear(Impl* queue, __code next(...));
        __code put(Impl* queue, union Data* data, __code next(...));
        __code take(Impl* queue, __code next(union Data*, ...));
        __code isEmpty(Impl* queue, __code next(...), __code whenEmpty(...));
} Queue;
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="interface--1">Interface の定義</h2>
<ul>
  <li><strong>__code next(…)</strong> は一種のクロージャであり、引数で Code Gear を渡すことが出来る</li>
  <li>クロージャの引数には Input Data Gear を指定する</li>
  <li>… は可変長引数のような扱いで、Code Gear が複数の Input Data Gear を待っている可能性がある</li>
</ul>

<pre lang="c"><code>typedef struct Queue&lt;Impl&gt;{
        // Data Gear parameter
        union Data* queue;
        union Data* data;
        __code next(...);
        __code whenEmpty(...);

        // Code Gear
        __code clear(Impl* queue, __code next(...));
        __code put(Impl* queue, union Data* data, __code next(...));
        __code take(Impl* queue, __code next(union Data*, ...));
        __code isEmpty(Impl* queue, __code next(...), __code whenEmpty(...));
} Queue;
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="interface--2">Interface の実装</h2>
<ul>
  <li>Interface には複数の実装を行うことが出来る</li>
  <li>Code Gear の番号を Interface の定義に代入することで実装を行う</li>
  <li>代入する Code Gear を入れ替えることで別の実装を表現する</li>
  <li>実装した Data Gear の生成は関数呼び出しで行われ、外から見るとInterface の型で扱われる</li>
</ul>

<pre lang="c"><code>Queue* createSingleLinkedQueue(struct Context* context) {
    struct Queue* queue = new Queue(); // Allocate Queue interface
    struct SingleLinkedQueue* singleLinkedQueue = new SingleLinkedQueue(); // Allocate Queue implement
    queue-&gt;queue = (union Data*)singleLinkedQueue;
    singleLinkedQueue-&gt;top  = new Element();
    singleLinkedQueue-&gt;last = singleLinkedQueue-&gt;top;
    queue-&gt;clear = C_clearSingleLinkedQueue;
    queue-&gt;put  = C_putSingleLinkedQueue;
    queue-&gt;take  = C_takeSingleLinkedQueue;
    queue-&gt;isEmpty = C_isEmptySingleLinkedQueue;
    return queue;
}
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<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の型で包んだData Gear、 <strong>method</strong> は実装した Code Gear に対応する</li>
</ul>

<pre lang="c"><code>__code code1() { 
    Queue* queue = createSingleLinkedQueue(context);
    Node* node = new Node();
    node-&gt;color = Red;
    goto queue-&gt;put(node, code2);
}
</code></pre>


</div>
<div class='slide '>
<!-- _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 の引数情報に合わせて引数に格納し、 実装された Code Gear へ継続する</li>
</ul>

<pre lang="c"><code>__code code1(struct Context *context) {
    Queue* queue = createSingleLinkedQueue(context);
    Node* node = &amp;ALLOCATE(context, Node)-&gt;Node;
    node-&gt;color = Red;
    Gearef(context, Queue)-&gt;queue = (union Data*) queue;
    Gearef(context, Queue)-&gt;data = (union Data*) node;
    Gearef(context, Queue)-&gt;next = C_code2;
    goto meta(context, queue-&gt;put);
}
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="interface--stub-code-gear">Interface での stub Code Gear</h2>
<ul>
  <li>メタ計算で格納された引数は stub Code Gear で Code Gear に渡される</li>
  <li>Interface を実装した Code Gear は Interface の定義から stub Code Gear の自動生成が可能</li>
</ul>

<pre lang="c"><code>// implement put code gear
__code putSingleLinkedQueue(struct Context *context,struct SingleLinkedQueue* queue,
                            union Data* data, enum Code next) {
    ...
}

// generated by script
__code putSingleLinkedQueue_stub(struct Context* context) {
	SingleLinkedQueue* queue = (SingleLinkedQueue*)GearImpl(context, Queue, queue);
	Data* data = Gearef(context, Queue)-&gt;data;
	enum Code next = Gearef(context, Queue)-&gt;next;
	goto putSingleLinkedQueue(context, queue, data, next);
} 
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-2">並列処理の構成</h2>
<ul>
  <li>今回は並列処理機構である
    <ul>
      <li>Task</li>
      <li>TaskManager
        <ul>
          <li>Worker の生成、依存関係を解決したTask を Worker に送信する</li>
        </ul>
      </li>
      <li>Worker
        <ul>
          <li>SynchronizedQueue から Task を一つずつ取得し、実行する</li>
          <li>Worker 毎に POSIX Therad などを生成し、それぞれのスレッドで Code Gear を実行する</li>
        </ul>
      </li>
      <li>SynchronizedQueue
        <ul>
          <li>マルチスレッド環境でもデータの同期が行われる Queue</li>
        </ul>
      </li>
    </ul>
  </li>
  <li>をInterface を用いて実装した</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="task">Task</h2>
<ul>
  <li>Gears OS では Context が並列実行の Task に相当する</li>
  <li>Context は Task用の情報として以下の情報をもっている
    <ul>
      <li>実行する Code Gear</li>
      <li>Input/Output Data Gear の格納場所</li>
      <li>待っている Input Data Gear の数</li>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="taskmanger">TaskManger</h2>
<ul>
  <li>Worker を作成、終了処理を行う
    <ul>
      <li>CPU、 GPU の数分生成する</li>
    </ul>
  </li>
  <li>依存関係を解決した Task を各 Worker の Queue に送信する</li>
</ul>

<div>
        <img src="./images/sendTask.svg" alt="message" style="float: left;width: 50%;" />
    <div style="float: left; width: 50%;">
        <ol>
            <li value="0">Task を Input Data Gear としてTaskManager の spawn を呼び出す</li>
            <li value="1">Task が待っている Data Gear のカウンタである IDGCount をチェックする</li>
            <li value="2">IDGCount が0の場合 Data Gear が 揃っているので Worker の Queue に Task を送信する</li>
        </ol>
    </div>
    <div style="clear: both;"></div>
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="worker">Worker</h2>
<ul>
  <li>初期化時に Worker 用の Context を生成し、Code Gear を実行していく</li>
  <li>TaskManager から送信された Task を一つずつ取得して実行する</li>
</ul>

<div>
    <img src="./images/workerRun.svg" alt="message" style="float: left;width: 50%;" />
    <div style="float: left; width: 50%;">
        <ol>
            <li>Worker は Queue から Task(Context)を取得する</li>
            <li>Worker の Context からTask の Context へ入れ替える</li>
                
            <li>Task に設定されている Code Gear を実行</li>
            <li>Task の Output Data Gear の書き出し</li>
            <li>Task Context から Worker の Context へ入れ替える</li>
            <li>Worker は再び Queue から Task を取得する</li>
        </ol>
    </div>
    <div style="clear: both;"></div>
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="synchronized-queue">Synchronized Queue</h2>
<ul>
  <li>Worker で使用される Queue は Task を送信するTaskManager と Task を取得する Worker 毎で操作される</li>
  <li>そのためマルチスレッドでのデータの同期処理を行える SynchronizedQueue として実装する</li>
  <li>Gears OS では CAS(Check and Set、 Compare and Swap) を使用した実装を行った
    <ul>
      <li>CAS は値を更新する際に更新前の値と実際に保存されているメモリ番地の値を比較し、変化がなければ値を更新する</li>
      <li>メモリ番地の値が変わっているなら、もう一度 CAS を行う</li>
    </ul>
  </li>
</ul>

<div style="text-align: center;">
    <img src="./images/synchronizedQueue.svg" alt="message" width="600" />
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-3">依存関係の解決</h2>
<ul>
  <li>依存関係の解決は Data Gear がメタレベルで持っている Queue を使用する</li>
  <li>この Queue には Data Gear に依存関係がある Context が格納されている</li>
</ul>

<div>
    <img src="./images/dependency.svg" alt="message" style="float: left;width: 50%;" />
    <div style="float: left; width: 50%;">
        <ol>
            <li>Task に設定されている Code Gear を実行する</li>
            <li>Output Data Gear の書き出し処理を行う際にメタレベルの Queue を参照する</li>
                
            <li>依存関係にある Task を取り出し、IDGCount をデクリメントする</li>
                
        </ol>
    </div>
    <div style="clear: both;"></div>
</div>

<ul>
  <li>IDGCountの値が0になった実行可能な Task は TaskManager を通して Worker に送信される</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-4">並列構文</h2>
<ul>
  <li>並列実行する際は新しく Context を生成し、実行する Code Gear、待ち合わせる Input Data Gear の数、Input/Output Data Gear への参照を設定する</li>
  <li>この記述を直接書くと Meta Data Gear である Context を直接参照しているため、ノーマルレベルでの記述としては好ましくない</li>
  <li>Context の設定は煩雑な記述だが、並列実行されることを除けば通常の CbC の goto 文と同等</li>
  <li>そこで Context を直接参照しない並列構文、 <strong>par goto</strong> 構文を新たに考案した</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="par-goto-">par goto 構文</h2>
<ul>
  <li>par goto 構文を記述すると新しく Context を生成し、TaskManager を通して Worker に送信される</li>
  <li>par goto 構文には引数として Input/Output Data Gear等を渡す</li>
  <li>スクリプトによって Code Gear の Input/Output の数を解析し、待っている Input Data Gear の数を設定する</li>
  <li>並列実行される Task は <strong>__exit</strong> に継続することで終了する
    <ul>
      <li>Gears OS の Task はOutput Data Gear を書き出す処理で終了するため <strong>__exit</strong> に直接継続せずに Data Gear を書き出す処理に継続する</li>
    </ul>
  </li>
  <li>par goto 構文は通常のプログラミングの関数呼び出しのように扱える</li>
</ul>

<pre lang="c"><code>__code code1(Integer *integer1, Integer * integer2, Integer *output) {
    par goto add(integer1, integer2, output, __exit);
    goto code2();
}
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="cuda-">CUDA への対応</h2>
<ul>
  <li>Gears OS は GPU での実行もサポートしている</li>
  <li>GPU で性能を出すためには GPU に合わせた並列プログラミングが必要になる</li>
  <li>CUDA は GPU を Device、 CPU を Host として定義する</li>
  <li>CUDA は処理の最小の単位を thread とし、それをまとめた block を展開し Device 上で実行されるプログラム(Kernel)を実行する</li>
  <li>今回、CUDA に合わせた並列処理機構である CUDAWorker、 CUDAExecutor をInterface を用いて実装し、 CPU、GPU間のデータのマッピングのために CUDABuffer を用意した</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="cudaworker">CUDAWorker</h2>
<ul>
  <li>CUDA で実行する Task を受け取る Worker</li>
  <li>初期化の際に CUDA ライブラリの初期化や CUDAExecutor の生成を行う</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="cudaexecutor">CUDAExecutor</h2>
<ul>
  <li>CUDAExecutor は Executor Interface を実装した Data Gear</li>
  <li>以下の Code Gear を実装している
    <ul>
      <li>HostからDevice へのデータの送信(read)</li>
      <li>kernel の実行(exec)</li>
      <li>Device から Host へのデータの書き出し(write)</li>
    </ul>
  </li>
</ul>

<pre lang="c"><code>typedef struct Executor&lt;Impl&gt;{
    union Data* Executor;
    struct Context* task;
    __code next(...);
    // method
    __code read(Impl* executor, struct Context* task, __code next(...));
    __code exec(Impl* executor, struct Context* task, __code next(...));
    __code write(Impl* executor, struct Context* task, __code next(...));
}
</code></pre>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="cudabuffer">CUDABuffer</h2>
<ul>
  <li>Host、Device 間でデータのやり取りをする際、 Gears OS での Data Gear をDevice 用にマッピングする必要がある</li>
  <li>CUDA Buffer ではそのマッピングを行う
    <ul>
      <li>このマッピングは Task に設定されている stub Code Gear で行われる</li>
    </ul>
  </li>
</ul>

<div style="text-align: center;">
    <img src="./images/cudaDataArchitecture.svg" alt="message" width="400" />
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="cuda--1">CUDA での呼び出し</h2>
<ul>
  <li>Gears OS では Task に設定している Code Gear の stub Code Gear で CUDA実行を切り替える</li>
  <li>CUDABuffer でのマッピング、実行する kernel の読み込みは stub Code Gear で行われる</li>
  <li>CUDA で実行する際は stub Code Gear に対応する Code Gear ではなく、 CUDAExecutor の Code Gear に継続する</li>
</ul>

<div style="text-align: center;">
    <img src="./images/gotoCUDAExecutor.svg" alt="message" width="600" />
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="gears-os-">Gears OS の評価</h2>
<ul>
  <li>並列処理のタスクの例題として Twice と BitonicSort を実装し、 以下の環境で測定を行った</li>
  <li>CPU 環境
    <ul>
      <li>Model : Dell PowerEdgeR630</li>
      <li>Memory : 768GB</li>
      <li>CPU : 2 x 18-Core Intel Xeon 2.30GHz</li>
    </ul>
  </li>
  <li>GPU 環境
    <ul>
      <li>GPU : GeForce GTX 1070</li>
      <li>Cores : 1920</li>
      <li>ClockSpeed : 1683MHz</li>
      <li>Memory Size : 8GB GDDR5</li>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="twice">Twice</h2>
<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 で kernel を実行</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="twice-">Twice の結果</h2>
<ul>
  <li>GPU は CPU との通信時間を含めた時間、 GPU(kernel only) は kernel のみの実行時間を示している</li>
  <li>1 CPU と 32 CPU では 約27.1倍の速度向上が見られた</li>
  <li>GPU は通信時間を含めると 8 CPU の約1.25倍、 kernel のみの実行では 32 CPU の約7.22倍になった</li>
  <li>通信時間がオーバーヘッドになっている</li>
</ul>

<table border="1" align="center" width="50%">
  <tbody>
    <tr>
      <td style="text-align: center;">Processor</td>
      <td style="text-align: center;">Time(ms)</td>
    </tr>
    <tr>
      <td style="text-align: center;">1 CPU</td>
      <td style="text-align: right;">1181.215</td>
    </tr>
    <tr>
      <td style="text-align: center;">2 CPUs</td>
      <td style="text-align: right;">627.914</td>
    </tr>
    <tr>
      <td style="text-align: center;">4 CPUs</td>
      <td style="text-align: right;">324.059</td>
    </tr>
    <tr>
      <td style="text-align: center;">8 CPUs</td>
      <td style="text-align: right;">159.932</td>
    </tr>
    <tr>
      <td style="text-align: center;">16 CPUs</td>
      <td style="text-align: right;">85.518</td>
    </tr>
    <tr>
      <td style="text-align: center;">32 CPUs</td>
      <td style="text-align: right;">43.496</td>
    </tr>
    <tr>
      <td style="text-align: center;">GPU</td>
      <td style="text-align: right;">127.018</td>
    </tr>
    <tr>
      <td style="text-align: center;">GPU(kernel only)</td>
      <td style="text-align: right;">6.018</td>
    </tr>
  </tbody>
</table>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="bitonicsort">BitonicSort</h2>
<ul>
  <li>並列処理向けのソートアルゴリズム</li>
  <li>決まった2点間の要素の入れ替えを並列に行うことでソーティングを進めていく</li>
  <li>要素数 2^24</li>
  <li>CPU での実行時は要素数 2^24 を 2^6 個に分割して Task を生成する</li>
  <li>GPU での実行時は1次元の block 数を 2^14、 block 内の thread 数を 2^10 で kernel を実行</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="bitonicsort-">BitonicSort の結果</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>
</ul>

<table border="1" align="center" width="50%">
  <tbody>
    <tr>
      <td style="text-align: center;">Processor</td>
      <td style="text-align: center;">Time(s)</td>
    </tr>
    <tr>
      <td style="text-align: center;">1 CPU</td>
      <td style="text-align: right;">41.416</td>
    </tr>
    <tr>
      <td style="text-align: center;">2 CPUs</td>
      <td style="text-align: right;">23.340</td>
    </tr>
    <tr>
      <td style="text-align: center;">4 CPUs</td>
      <td style="text-align: right;">11.952</td>
    </tr>
    <tr>
      <td style="text-align: center;">8 CPUs</td>
      <td style="text-align: right;">6.320</td>
    </tr>
    <tr>
      <td style="text-align: center;">16 CPUs</td>
      <td style="text-align: right;">3.336</td>
    </tr>
    <tr>
      <td style="text-align: center;">32 CPUs</td>
      <td style="text-align: right;">1.872</td>
    </tr>
    <tr>
      <td style="text-align: center;">GPU</td>
      <td style="text-align: right;">5.420</td>
    </tr>
    <tr>
      <td style="text-align: center;">GPU(kernel only)</td>
      <td style="text-align: right;">0.163</td>
    </tr>
  </tbody>
</table>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="openmp-">OpenMP との比較</h2>
<ul>
  <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="go-">Go 言語との比較</h2>
<ul>
  <li>Go 言語は並列実行を <strong>go funciton(argv)</strong> の構文で行う。 この実行を goroutine と呼ぶ</li>
  <li>goroutine 間のデータの待ち合わせはチャネルというデータ構造で行う</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="section-5">まとめ</h2>
<ul>
  <li>Gears OS の並列処理機構を Interface を用いて実装を行った</li>
  <li>Interface を導入することで、見通しの良し Gears OS のプログラミングが可能となった</li>
  <li>par goto  構文を導入することで、ノーマルレベルで並列処理の記述が可能になった</li>
  <li>2つの例題である程度の台数効果が出ることを確認した</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-6">今後の課題</h2>
<ul>
  <li>Gears OS の並列処理の信頼性の保証、チューニングを行う</li>
  <li>Gears OS では証明とモデル検査をメタレベルで実現することで信頼性を保証する
    <ul>
      <li>証明は CbC のプログラムを証明支援系の Agda に対応して行う。 並列処理の信頼性を保証するには SynchronizedQueue の証明を行う必要がある</li>
      <li>モデル検査は CbC で記述された モデル検査器である akasha を使用して行う。 モデル検査の方針としては Code Gear の並列実行を擬似並列で実行し、全ての組合せを列挙する方法で行う</li>
    </ul>
  </li>
  <li>現在の CUDA 実装では CPU、GPU 間のデータの通信コストがかかってしまうことが例題からわかった
    <ul>
      <li>Meta Data Gear に Data Gear が CPU、 GPU のどこで所持されているのかを持たせ、 GPU の Data Gear が CPU で必要になったときに始めてデータの通信を行う</li>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-7">今後の課題</h2>
<ul>
  <li>OpenMP、 Go 言語で Twice を実装し、 Gears OS の性能比較を行った</li>
  <li>その結果、 Gears OS が 1CPU での動作が遅いということがわかった。
    <ul>
      <li>par goto 文を使用する度に Context を生成するため、 ある程度の時間がかかってしまう</li>
      <li>モデル検査で par goto で実行する Code Gear のフローを解析し、処理が軽い場合は Context を生成せずに関数呼出しを行う等の最適化が必要</li>
    </ul>
  </li>
</ul>

<div style="text-align: center;">
    <img src="./images/compareExamples.svg" alt="message" width="500" />
</div>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-8">データ並列</h2>
<ul>
  <li>data並列はあるデータ構造がサブデータへ分割することが可能で、各サブデータに行う処理が同じ場合に有効な並列処理手法</li>
  <li>Gears OS ではdata 並列は par goto 構文に<strong>iterate(分割数)</strong>を追加することで可能になる</li>
  <li>データ並列の Task は CPU で実行する際は Task にインデックスを付与して分割数分コピーして実行する</li>
  <li>CUDA の場合は Kernel を実行する際にパラメーターとして分割数を渡す</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="task-">Task 間の同期処理</h2>
<ul>
  <li>Context 間での同期処理を行うために Semaphore を実装</li>
  <li>Semaphore はContext 停止用の待ち Queue を持つ</li>
</ul>

<div style="text-align: center;">
    <img src="./images/semaphoreSequence.svg" alt="message" width="700" />
</div>

<!--
## OpenMP との比較
- 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 言語でも 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>
-->
<!-- === end markdown block === -->
</div>


</div><!-- presentation -->
</body>
</html>