Mercurial > hg > Papers > 2018 > parusu-master

<!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-12 14:45:13 +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>メタ計算 の処理も Code/Data 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="context">Context</h2>
<ul>
  <li>Context は接続可能な Code/Data Gear の集合を表現する Meta Data Gear</li>
  <li>従来のOS のスレッドやプロセスに対応し、以下の要素を定義している
    <ul>
      <li>独立したメモリ空間</li>
      <li>Code/Data Gear へのポインタ</li>
      <li>並列実行用の Task 情報</li>
      <li>Data Gear の型情報</li>
    </ul>
  </li>
  <li>Gears OS ではメタ計算でこの Context を経由して Data Gear にアクセスする</li>
</ul>


</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="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>// 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>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="context--stub-code-gear-">Context での stub Code Gear の記述の問題点</h2>
<ul>
  <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>stub Code Gear を自動生成するために、同じ番号の Data Gear  を使いまわす問題</li>
    </ul>
  </li>
  <li>そのため Gears OS のモジュール化する仕組みとして <strong>Interface</strong> を導入した</li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="interface">Interface</h2>
<ul>
  <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>引数のData Gear 群</li>
      <li>操作(API) 実行後に継続される Code Gear</li>
      <li>操作(API) である Code Gear と 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>Interface には複数の実装を行うことが出来る</li>
  <li>実装した Code Gear を Interface で定義した Code Gear に代入することで実装を行う</li>
  <li>代入する Code Gear を入れ替えることで別の実装を表現する</li>
  <li>実装した Data Gear の生成は関数呼び出しで行われ、外から見るとInterface の型で扱われる</li>
</ul>

<pre><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--2">Interface の実装例</h2>
<ul>
  <li>SingleLinkedQueue の put 実装</li>
  <li>引数は Queue Interface の put 定義にあわせる</li>
  <li>第1引数は 実装対象の Data Gear の型になる</li>
  <li>第3引数の(…) は Output Data Gear を記述する
    <ul>
      <li>… は可変長引数のような扱いで、 継続先の Code Gear が複数の値をInput Data Gear とする可能性がある</li>
    </ul>
  </li>
</ul>

<pre lang="c"><code>__code putSingleLinkedQueue(struct SingleLinkedQueue* queue, union Data* data, __code next(...)) {
    Element* element = new Element();
    element-&gt;data = data;
    element-&gt;next = NULL;
    queue-&gt;last-&gt;next  = element;
    queue-&gt;last = element;
    goto next(...);
}
</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><code>__code code1() {
    Queue* queue = createSingleLinkedQueue(context);
    Node* node = new Node();
    node-&gt;color = Red;
    goto queue-&gt;put(node, queueTest2);
}
</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><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_queueTest2;
    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>__code putSingleLinkedQueue(struct Context *context,struct SingleLinkedQueue* queue, union Data* data, enum Code next) {
    Element* element = &amp;ALLOCATE(context, Element)-&gt;Element;
    element-&gt;data = data;
    element-&gt;next = NULL;
    queue-&gt;last-&gt;next  = element;
    queue-&gt;last = element;
    goto meta(context, 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>今回はInterface を利用した並列処理機構の実装を行う</li>
  <li>構成要素として以下が挙げられる
    <ul>
      <li>Task(Context)</li>
      <li>TaskManager
        <ul>
          <li>Worker の生成、依存関係を解決したTask を Worker に送信する</li>
        </ul>
      </li>
      <li>Worker
        <ul>
          <li>SynchronizedQueue から Task を取得し、実行する</li>
        </ul>
      </li>
      <li>SynchronizedQueue
        <ul>
          <li>マルチスレッド 環境でもデータの一貫性を保証する Queue</li>
        </ul>
      </li>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="taskcontext">Task(Context)</h2>
<ul>
  <li>Gears OS では並列実行する Task を Context で表現する</li>
  <li>Context は Task用の情報として以下の情報をもっている
    <ul>
      <li>実行する Code Gear</li>
      <li>Input/Output Data Gear の格納場所</li>
      <li>待っている Input Data Gear の数</li>
    </ul>
  </li>
  <li>実際に実行される Code Gear の引数情報は Interface の Code Gear 実装と同等に記述できる
    <ul>
      <li>Code Gear の stub Code Gear は自動生成される</li>
    </ul>
  </li>
</ul>

<pre lang="c"><code>__code add(struct Integer* input1, struct Integer* input2, __code next(struct Integer* output, ...)) {
    output-&gt;value = input1-&gt;value + input2-&gt;value;
    goto next(output, ...);
}
</code></pre>


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

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


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="worker">Worker</h2>
<ul>
  <li>初期化の際にスレッドと Worker 用の Context を生成する</li>
  <li>生成されたスレッドではTaskManager から送信された Task を取得して実行する</li>
</ul>

<div>
    <div style="float: left;">
        <img src="./images/workerRun.svg" alt="message" width="600" />
    </div>
    <div style="float: left; font-size=100%;">
        <ol>
            <li>Worker は Queue から Task を取得する</li>
            <li>Worker Context から Task へ入れ替える</li>
            <li>Task の Code Gear を実行</li>
            <li>Task の Output Data Gear の書き出し</li>
            <li>Task から WorkerContext へ入れ替える</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</li>
  <li>Task を送信するスレッドと Task を取得するスレッドで操作される</li>
  <li>そのためマルチスレッドでのデータの一貫性を保証する必要がある</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>

<pre><code>struct SynchronizedQueue {
    struct Element* top;
    struct Element* last;
    struct Atomic* atomic;
};
// Singly Linked List element
struct Element {
    union Data* top;
    struct Element* next;
};
</code></pre>


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

<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 '>
<!-- _S9SLIDE_ -->
<h2 id="section-4">並列構文</h2>
<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>そこで Context を直接参照しない並列構文、 <strong>par goto</strong> 構文を新たに考案した</li>
  <li>par goto 構文には引数として Input/Output Data Gear等を渡す
    <ul>
      <li>スクリプトによって Code Gear の Input/Output の数を解析する</li>
    </ul>
  </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 への実行のサポートをおこなった</li>
  <li>CUDA は GPU を Device、 CPU を Host として定義する</li>
  <li>CUDA は処理の最小の単位を thread とし、それをまとめた block を展開し Device 上で実行されるプログラム(Kernel)を実行する</li>
  <li>今回 CUDAWorker、CUDAExecutor、 CUDABuffer を使用して CUDA に合わせた並列処理機構を提供する</li>
</ul>


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


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="cudaexecutor">CUDAExecutor</h2>
<ul>
  <li>CUDAExecutor は Executor Interface を実装した以下の 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 用にマッピングする必要がある
    <ul>
      <li>Device にデータ領域を確保するにはサイズの指定が必要</li>
      <li>Data Gear には Meta Data Gear でデータのサイズを持っている</li>
      <li>しかし、 Data Gear の要素の中に Data Gear へのポインタがあるとポインタ分でサイズ計算してしまうため、 GPU では参照できなくなってしまう</li>
    </ul>
  </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="500" />
</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>stub Code Gear で CUDABuffer でのマッピング、 実行する kernel の読み込みを行う</li>
  <li>stub Code Gear は CUDA で実行する際は CUDAExecutor の Code Gear に継続する</li>
</ul>


</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>
      <li>Memory Bandwidth :  256GB/s</li>
    </ul>
  </li>
</ul>


</div>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="twice">Twice</h2>
<ul>
  <li>Twice は与えられた整数配列を2倍にする例題である</li>
  <li>並列実行の依存関係がなく、並列度が高い課題である</li>
</ul>


</div>
<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>
      <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点間の要素の入れ替えをステージ毎に並列に実行し、 Output Data Gear として書き出し、次のステージの Code Gear の Input Data Gear とする</li>
</ul>

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


</div>
<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>
    <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="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="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>
<div class='slide '>
<!-- _S9SLIDE_ -->
<h2 id="section-5">まとめ</h2>
<ul>
  <li>Gears OS の並列処理機構の実装を行った</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>OpenMP、 Goとの比較から、 Gears OS が 1CPU での動作が遅いということがわかった。
    <ul>
      <li>par goto 文を使用する度に Context を生成するため、 ある程度の時間がかかってしまう</li>
      <li>モデル検査で par goto の Code Gear のフローを解析し、処理がかる場合は Context を生成せずに関数呼出しを行う等の最適化が必要</li>
    </ul>
  </li>
  <li>現在の CUDA 実装では CPU、GPU 間のデータの通信コストがかかってしまうことが例題からわかった
    <ul>
      <li>Meta Data Gear に Data Gear が CPU、 GPU のどこで所持されているのかを持たせ、 GPU の Data Gear が CPU で必要になったときに始めてデータの通信を行う</li>
    </ul>
  </li>
</ul>
<!-- === end markdown block === -->
</div>


</div><!-- presentation -->
</body>
</html>
author	Tatsuki IHA <innparusu@cr.ie.u-ryukyu.ac.jp>
date	Mon, 12 Feb 2018 14:48:30 +0900
parents	3c127f675c45
children	a9885c038bb6