# HG changeset patch
# User Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp>
# Date 1424193050 -32400
# Node ID 4c245ed4e61a3ad12614244ddc4049579ebc0c9e
# Parent  3a35d13818e5548d92cdfff52f43337bb2d68bfb
io thread slide

diff -r 3a35d13818e5 -r 4c245ed4e61a paper/images/blockedread.pdf
Binary file paper/images/blockedread.pdf has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a paper/images/wordcount.pdf
Binary file paper/images/wordcount.pdf has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a paper/master_paper.pdf
Binary file paper/master_paper.pdf has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a slide/blank.html
--- a/slide/blank.html	Wed Feb 18 00:15:07 2015 +0900
+++ b/slide/blank.html	Wed Feb 18 02:10:50 2015 +0900
@@ -474,7 +474,7 @@
           これらのフレームワークを用いて Cerium に GPU 上で 並列実行する機能を加えた。
         </p>
         <p>
-          Scheduler から受け取った Task やデータをOpenCL、CUDA の API を介して GPU に転送する機構、
+          TaskManager から受け取った Task やデータをOpenCL、CUDA の API を介して GPU に転送する機構、
           GpuScheduler と CudaScheduler を実装した。
         </p>
         <div align="center">
@@ -484,30 +484,189 @@
       
       <div class='slide'>
         <h2>フレームワークを用いた GPU の制御</h2>
+        <p>
+          GpuScheduler、CudaScheduler ではそれぞれのフレームワークを用いて GPU の制御を行っている。
+          行われていることは以下の3つに分けられる。
+        </p>
+        <ul>
+          <li>Host から Device へのデータ転送
+          <li>kernel の実行
+          <li>Device から Host へのデータ転送
+        </ul>
+        <p>
+           CommandQueue と呼ばれる機構を用いてこういった GPU を制御するための処理を行っていく。
+           CommandQueue に命令を起こるためのしくみで、制御は全てこの Queue を介して行われる。
+        </p>
+        <p>これらはRead, Exec、Write に対応する。
+          GPGPU 用の Scheduler でもパイプラインを構成する。</p>
       </div>
 
       <div class='slide'>
-        <h2>GPGPU におけるパイプラインの実装</h2>
+        <h2>GPGPU におけるパイプラインの実装(Read)</h2>
+        <p>
+          GpuScheduler では SchedTask を用いてない。
+          メインループでは2つの CommandQueue を保持し、GPU の制御命令を二段のパイプラインで実行していく。
+          TaskList から Task を取り出し、Task から実行する kernel やパラメタを生成し、
+          各種フレームワークの API を通して GPU のメモリに転送される。
+        </p>
+
+        <p>
+          全ての Task が終了すると SynchronizedQueue を通してTaskManager に終了を通知する。
+        </p>
+
+        <p>
+          Scheduler の内部で Platform や DeviceID の取得、
+          kernel の build や load といった API を使用するための初期化も行っており、
+          並列化したい処理のみに集中できる。
+        </p>
       </div>
 
       <div class='slide'>
         <h2>GPGPU におけるデータ並列</h2>
+        <p>
+          マルチコア CPU と同様に、GPGPU に関してもデータ並列実行をサポートした。
+        </p>
+        <p>
+          GPU 上でデータ並列実行する際も iterate API によりデータ並列用の Task を生成できる。
+          生成部分の記述はマルチコア CPU と同じ形式で記述できる。
+          また、Task 自体の記述もほぼ同じ形式となる。以下に Task の例を示す。
+        </p>
+
+        <table>
+          <tr align="left">
+            <th>
+              <pre class="code">
+__kernel void // OpenCL
+multiply(__global const long  *params,
+         __global const float *input1,
+         __global const float *input2,
+         __global const float *output) {
+
+    long id = get_global_id(0);
+
+    output[id] = input1[id] * input2[id];
+}
+              </pre>
+            </th>
+            <th>
+              <pre class="code">
+__global__ void // CUDA
+multiply(__global const long  *params,
+         __global const float *input1,
+         __global const float *input2,
+         __global const float *output) {
+
+    int id = blockIdx.x * blockDim.x + threadIdx.x;
+
+    output[id] = input1[id] * input2[id];
+}
+              </pre>
+            </th>
+          </tr>
+        </table>
+      </div>
+
+      <div class='slide'>
+        <h2>並列処理向け I/O</h2>
+        <p>
+          ファイルの読み込みなどの I/O を含むプログラムは、
+          読み込み時間が Task のと比較してオーバーヘッドになることが多い。
+          プログラムの並列化を行ったとしても I/O がボトルネックになってしまうと処理は高速にならない。
+        </p>
+        <p>並列計算と同時に動作する、並列 I/O の実装を行った。</p>
       </div>
 
       <div class='slide'>
         <h2>Cerium の I/O(mmap による読み込み)</h2>
+        <p>
+          Cerium ではファイルの読み込みを mmap で行っていた。</p>
+          <ul>
+            <li>mmap はまず仮想メモリにファイルをマッピングする。
+            <li>マッピングしたメモリ空間にアクセスがあったら OS が読み込みを行う。
+            <li>mmap は並列に動作せず、逐次処理
+            <li>読み込みが OS 依存となり、環境に左右されやすい
+        <p>並列に動作する I/O の機構が必要である</p>
+      </div>
+
+      
+      <div class='slide'>
+        <h2>WordCount</h2>
+        <p>サイズの大きいファイルを読み込む例題、WordCount を元に並列 I/O について考える。</p>
+        <p>
+          WordCount は Input としてファイルを受け取り、ファイルの単語数と行数を集計して表示する例題である。
+        </p>
+        <table>
+          <tr><th><img src="./images/wordcount.png" width="600"></th>
+            <th align="left">
+              <ul>
+                <li>input ファイルを一定の大きさ分割する
+                <li>読み込んだテキストファイルに対してそれぞれ並列に計算を行う
+                <li>PrintTask が計算結果を集計して出力する
+              </ul>
+            </th>
+          </tr>
+        </table>
+      </div>
+      
+      <div class='slide'>
+        <h2>BlockedRead による I/O の並列化</h2>
+        <p>ファイルを読み込んで、読み込んだファイルに対して並列実行を行う場合、ファイルを分割して処理を行う。</p>
+        <p>よって読み込みの処理自体を分割し、ある程度の大きさ(Block)ごとに読み込みと Task の実行を行う。</p>
+        <p>読み込みの処理自体を分割して行う。これを BlockedRead と呼ぶ。</p>
+        </p>
       </div>
 
       <div class='slide'>
-        <h2>BlockedRead による I/O の並列化</h2>
+        <h2>BlockedRead を用いた WordCount</h2>
+        <div align="center">
+          <img src="./images/blockedread.png" width="600">
+        </div>
+        <p>
+          BlockedRead を用いて WordCount を行う際、読み込み用の Task と
+          読み込んだファイルに対して処理を行う Task の2つを生成する。
+        </p>
+        <p>ファイルを分割して読み込み、
+          読み込んだファイルに対して WordCount を行う一定数のTask(BlockedTask)を割り当てる。
+          Task には依存関係を設定する必要があり、図のTask n+1 はTask nを待つ必要がある。
+        </p>
+        <p>まだ読み込みが終了していない領域に割り当てられた Task が起動してしまう事を防ぐためである。</p>
+        <p>この wait によるロックはオーバーヘッドとなるため、なるべく発生しないことが望ましい。</p>
       </div>
 
       <div class='slide'>
         <h2>I/O 専用のThread</h2>
+        <p>
+          BlockedRead の依存関係による wait はなるべく発生しないことが望ましい。
+          そのため、BlockedRead は連続で Task の起動を行う必要がある。
+        </p>
+        <p>
+          Cerium には SPE_ANY という Thread があり、この Thread で Task の実行を行うと自動で実行するコアを割り振る。
+          しかし、SPE_ANY で BlockedRead を実行すると BlockedRead 間に別の Task が割り込んでしまう場合がある。
+        </p>
+        <div align="center">
+          <img src="./images/speblockedread.png" width="700">
+        </div>
+        <p>TaskBlock の依存関係によっては wait がかかってしまう。そこで、I/O 専用の Thread を作成した。</p>
       </div>
 
       <div class='slide'>
-        <h2>実験に利用する例題-Sort-</h2>
+        <h2>I/O 専用のThread</h2>
+        <p>
+          IO 専用の Thread を作成したが、それだけでは問題は解決しない場合がある。
+          IO thread 内では割り込みが生じる可能性はないが、thread レベルで割り込みが起きる可能性がある。
+          IO thread-SPE_ANY-IO Thread のような実行順序となる場合である。
+        </p>
+        <div align="center">
+          <img src="./images/iothread.png" width="700">
+        </div>
+        <p>
+          そのため、pthread_getschedparam() という POSIX スレッドの API を用いて IO Thread の priority を高く設定した。
+          IO Thread は必ず連続で行われることになる。
+        </p>
+      </div>
+
+      <div class='slide'>
+        <h2>-</h2>
       </div>
 
       <div class='slide'>
diff -r 3a35d13818e5 -r 4c245ed4e61a slide/images/blockedread.png
Binary file slide/images/blockedread.png has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a slide/images/iothread.png
Binary file slide/images/iothread.png has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a slide/images/speblockedread.pdf
Binary file slide/images/speblockedread.pdf has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a slide/images/speblockedread.png
Binary file slide/images/speblockedread.png has changed
diff -r 3a35d13818e5 -r 4c245ed4e61a slide/images/wordcount.png
Binary file slide/images/wordcount.png has changed