Mercurial > hg > Papers > 2014 > kkb-sigos

changeset 8:a5fc2e56aef7

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
revision
author Shohei KOKUBO <e105744@ie.u-ryukyu.ac.jp>
date Tue, 13 May 2014 15:45:05 +0900
parents 18b7450b6959
children 9bb5fca606b3
files presen/s6/images/ndrange_arch.png presen/s6/index.html
diffstat 2 files changed, 105 insertions(+), 61 deletions(-) [+]
line wrap: on
line diff
Binary file presen/s6/images/ndrange_arch.png has changed
--- a/presen/s6/index.html	Mon May 12 03:37:44 2014 +0900
+++ b/presen/s6/index.html	Tue May 13 15:45:05 2014 +0900
@@ -108,6 +108,11 @@
                 <hr style="color:#ffcc00;background-color:#ffcc00;text-align:left;border:none;width:300%;height:0.2em;">
             </div></td>
           </tr>
+          <tr>
+            <td><div align="left" style="margin-top:-5%">
+                担当教官：河野真治
+            </div></td>
+          </tr>
         </table>
       </div>
       
@@ -127,16 +132,16 @@
       <div class='slide'>
         <h2>はじめに(2/2)</h2>
         <p>
-          GPU や Cell などメモリ空間が異なるプロセッサはデータ転送がオーバーヘッドとなるので、データ転送を効率的に行えるかどうかで処理時間が大きく変わる。
+          GPU や Cell などメモリ空間が異なるプロセッサはデータ転送がオーバーヘッドとなる。このオーバーヘッドを解消するためには<font color="red">パイプライン処理</font>を行い、<font color="red">データ転送をオーバーラップ</font>する必要がある。
         </p>
         <p>
-          今回、Cerium に OpenCL, CUDA を用いた Scheduler を実装した。
+          様々なプロセッサを統合して扱えるフレームワークとして INRIA が開発している StarPU がある。StarPU は Cerium と同じタスクベースの非同期処理を行うフレームワークである。StarPU はパイプライン処理をサポートしていないので、パイプライン処理を行いたい場合は自分で実装するしかない。
         </p>
         <p>
-          Scheduler は<font color="red">自動でデータ転送をオーバーラップ</font>し、<font color="red">パイプライン処理</font>を行うように設計した。
+          今回、Ceirum に OpenCL, CUDA を用いた Scheduler を実装した。
         </p>
         <p>
-          
+          OpenCL, CUDA によるすべての最適化を Cerium では自動で行う。その一環として、実装した Scheduler はパイプライン処理によるデータ転送の効率化を自動で行うようになっている。
         </p>
       </div>
 
@@ -151,26 +156,7 @@
                   CPU と GPU はメモリ空間が異なるのでメモリの共有ができない。データにアクセスするためには空間ごとコピーする必要がある。この転送がオーバーヘッドとなる。
                 </p>
                 <p style="font-size:100%">
-                  つまり、GPGPU では<font color="red">データ転送を頻繁に行う</font> Task は並列度が低くなる。
-                </p>
-              </td>
-            </tr>
-          </tbody>
-        </table>
-      </div>
-      
-      <div class='slide'>
-        <h2>データ転送</h2>
-        <table border="0" cellpadding="0" cellspacing="0">
-          <tbody>
-            <tr>
-              <td><img src="images/gpu_data_parallel.png" style="height:350px"></td>
-              <td>
-                <p style="font-size:100%">
-                  タスク並列で実行する場合、タスクが担当するデータを実行する数だけデータ転送を行う。
-                </p>
-                <p style="font-size:100%">
-                  データ並列で実行する場合、すべてのデータを一度で転送する。kernel で担当する範囲を計算し、その部分に対して処理を行う。データ転送の回数が少なくなるので、OpenCL, CUDA ではデータ並列による実行が推奨されている。
+                  つまり、GPGPU では<font color="red">データ転送を頻繁に行う</font>タスクは並列度が低くなる。
                 </p>
               </td>
             </tr>
@@ -179,15 +165,34 @@
       </div>
       
       <div class='slide'>
+        <h2>並列化手法とデータ転送</h2>
+        <p>
+          並列化手法は大まかにタスク並列とデータ並列に分けることができる。
+        </p>
+        <p>
+          タスク並列では複数のプロセッサで異なる演算を行い、データ並列では複数のプロセッサで同じ演算を行う。
+        </p>
+        <p>
+          タスク並列では、同じ演算を行う場合でもタスクの数だけデータ転送を行う必要がある。
+        </p>
+        <p>
+          データ並列では、すべてのデータを一度で転送する。タスク側で担当する範囲を計算し、その部分に対して処理を行う。
+        </p>
+        <p>
+          同じ演算を行いたい場合、データ転送が少なくなるのでデータ並列の方が有効である。特に GPGPU ではデータ転送がオーバーヘッドとなるのでその差が顕著に現れる。
+        </p>
+      </div>
+      
+      <div class='slide'>
         <h2>OpenCL, CUDA による GPGPU への対応</h2>
         <p>
-          OpenCL では CommandQueue、CUDA では Stream をそれぞれ用いて GPU に命令を転送する。
+          OpenCL では CommandQueue、CUDA では Stream という同期機構がある。それぞれに命令を投入することで GPU を制御する。
         </p>
         <p>
           CommandQueue, Stream に投入された命令は発行された順序で実行される。これらを複数用いることで<font color="red">命令を並列に実行する</font>ことが可能になる。
         </p>
         <p>
-          OpenCL および CUDA では Task を kernel という単位で記述する。
+          OpenCL および CUDA ではタスクを kernel という単位で記述する。
         </p>
         <p>
           Cerium において、CPU/Cell 上で実行する場合と GPU 上で実行する場合、それぞれ<font color="red">ほぼ同じ形式の記述</font>が可能である。
@@ -195,20 +200,78 @@
       </div>
 
       <div class='slide'>
-        <h2>Multi Stream による並列実行</h2>
+        <h2>複数の同期機構を用いた並列実行</h2>
         <p>
-          Stream に投入された命令は投入された順序で実行されることが保証される。
+          CommandQueue, Stream に投入された命令は投入された順序で実行されることが保証される。
         </p>
         <p>
-          異なる Stream の命令の依存関係が解消され実行可能な状態の場合、命令は並列に実行される。これは OpenCL で CommandQueue を用いて実行する場合も同様である。
+          異なる CommandQueue, Stream の命令の依存関係が解消され実行可能な状態の場合、命令は並列に実行される。
+        </p>
+        <img src="images/parallel_exec.png" style="height:190px">
+        <p>
+          HtoD, kernel, DtoH は cuMemcpyHtoD, cuLaunchKernel, cuMemcpyDtoH という CUDA の API である。ホストからデバイスへのデータ転送、kernel の実行、デバイスからホストへのデータ転送をそれぞれ行う。OpenCL では clEnqueueReadBuffer, clEnqueueTask, clEnqueueWriteBuffer に相当する。
+        </p>
+      </div>
+
+      <div class='slide'>
+        <h2>命令の投入方法</h2>
+        <p>
+          注意点として CUDA では以下のように命令を投入しないと命令が並列に実行されない。
         </p>
-        <img src="images/parallel_exec.png" style="width:800px">
+        <pre style="font-size:100%">
+// BAD
+for (int i=0;i&lt;stream_num;i++) {
+    cuMemcpyHtoDAsync();
+    cuLaunchKernel();
+    cuMemcpyDtoHAsync();
+}</pre>
+<pre class="code" style="font-size:100%">
+// GOOD
+for (int i=0;i&lt;stream_num;i++)
+    cuMemcpyHtoDAsync();
+             
+for (int i=0;i&lt;stream_num;i++)
+    cuLaunchKernel();
+
+for (int i=0;i&lt;stream_num;i++)
+    cuMemcpyDtoHAsync();</pre>
+      </div>
+
+      <div class='slide'>
+        <h2>host の記述</h2>
+        <p>
+          2つの input の積を取り、output に返す例題の host
+        </p>
+        <pre style="font-size:90%">
+void
+multi_init(TaskManager *manager)
+{
+    A = new float[length];
+    B = new float[length];
+    C = new float[length];
+    for(int i=0; i&lt;length; i++) {
+        A[i]=(float)(i+1000);
+        B[i]=(float)(i+1)/10.f;
+    }
+    
+    HTask* multiply = manager-&gt;create_task(MULTIPLY_TASK);
+    multiply-&gt;set_cpu(CPU_TYPE);
+
+    multiply-&gt;set_inData(0,(memaddr)A, sizeof(float)*length);
+    multiply-&gt;set_inData(1,(memaddr)B, sizeof(float)*length);
+
+    multiply-&gt;set_outData(0,(memaddr)C, sizeof(float)*length);
+
+    multiply-&gt;iterate(length); 
+}</pre>
+        <p>
+          CPU, GPU ともに完全に同じ記述で Task を生成することができる。set_cpu で実行するデバイスを選択することができる。
       </div>
 
       <div class='slide'>
         <h2>kernel の記述</h2>
         <p>
-          2つの input の積を取り、output に返す例題 Multiply 
+          2つの input の積を取り、output に返す例題の kernel
         </p>
         <pre style="font-size:70%">
 // MultiCore CPU
@@ -253,13 +316,13 @@
               <td><img src="images/ndrange_arch.png" style="height:350px"></td>
               <td>
                 <p style="font-size:100%">
-                  3D グラフィックのような多次元のデータを処理する場合に高い並列度を保つには、データを分割して並列に実行する機能が必要である。
+                  1つの記述から ID の異なる複数の kernel を自動生成する。
                 </p>
                 <p style="font-size:100%">
-                  これを OpenCL, CUDA ではデータ並列と呼んでいる。
+                  OpenCL では get_global_id 関数、CUDA では組み込み変数を参照することで ID を取得することができる。
                 </p>
                 <p style="font-size:100%">
-                  1つの記述から ID の異なる複数の kernel を自動生成する。OpenCL では get_global_id 関数、CUDA では組み込み変数の blockDim, blockIdx, threadIdx を参照することで index を取得することができ、index から担当する範囲を決定し処理を行う。
+                  取得した ID から担当する範囲を計算し、その部分に対して処理を行う。
                 </p>
               </td>
             </tr>
@@ -312,21 +375,23 @@
       </div>
 
       <div class='slide'>
-        <h2>Cerium におけるデータ並列のための API(1/2)</h2>
+        <h2>CPU におけるデータ並列のための API</h2>
         <table border="0" cellpadding="0" cellspacing="0">
           <tbody>
             <tr>
               <td><img src="images/iterate.png" style="height:300px"></td>
               <td>
                 <p style="font-size:100%">
-                  Task に ID と input/output data を割り当てる「iterate」という API を実装した。
+                  データ並列で実行する API として「iterate」を実装した。
                 </p>
                 <ul>
-                  <li style="font-size:100%">iterate</li>
+                  <li style="font-size:100%">CPU におけるデータ並列</li>
                   <ul>
-                    <li style="font-size:100%">1つの記述から複数の Task を生成する。</li>
-                    <li style="font-size:100%">生成した複数の Task に ID と input/output data を割り当てる。</li>
-                    <li style="font-size:100%">Multi Core CPU でデータ並列を行う場合は SchedTask に index が設定されており、それを参照することで担当する範囲を決定し処理を行う。</li>
+                    <li style="font-size:100%">Task を複製する。</li>
+                    <li style="font-size:100%">複製した Task に index と input/output data を割り当て、各 CPU に転送する。</li>
+                    <li style="font-size:100%">各 CPU にも ID を割り当てる。</li>
+                    <li style="font-size:100%">各 CPU に割り当てられた ID から実行する Task を選択する。</li>
+                    <li style="font-size:100%">Task 側では SchedTask を参照することで index を取得することができる。</li>
                   </ul>
                 </ul>
               </td>
@@ -336,27 +401,6 @@
       </div>
 
       <div class='slide'>
-        <h2>Cerium におけるデータ並列のための API(2/2)</h2>
-<pre class="code" style="font-size:90%">
-void
-HTask::iterate(long x, long y, long z) {
-    this->flag.dim_count = 1;
-    TaskList *tl = (TaskList*)rbuf;
-    for (;tl;tl=tl->next) {
-        tl->set_last(last);
-        tl->dim=1;
-        tl->x=x;
-        tl->y=y;
-        tl->z=z;
-    }
-    mimpl->spawn_task(this);
-}</pre>
-        <p>
-          iterate を使用するとデータ並列で実行する Task として登録される。この時点で TaskList にパラメータが設定され、TaskManager から Scheduler に転送される。
-        </p>
-      </div>
-
-      <div class='slide'>
         <h2>ベンチマーク</h2>
         <ul>
           <li>測定環境</li>