Mercurial > hg > Papers > 2015 > yuhi-master
changeset 64:4c245ed4e61a
io thread slide
| author | Yuhi TOMARI <yuhi@cr.ie.u-ryukyu.ac.jp> |
|---|---|
| date | Wed, 18 Feb 2015 02:10:50 +0900 |
| parents | 3a35d13818e5 |
| children | a41beec3553e |
| files | paper/images/blockedread.pdf paper/images/wordcount.pdf paper/master_paper.pdf slide/blank.html slide/images/blockedread.png slide/images/iothread.png slide/images/speblockedread.pdf slide/images/speblockedread.png slide/images/wordcount.png |
| diffstat | 9 files changed, 163 insertions(+), 4 deletions(-) [+] |
line wrap: on
line diff
--- 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'>